This is the mail archive of the fortran@gcc.gnu.org mailing list for the GNU Fortran project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [gomp4] OpenACC wait directive


On 10/06/2014 03:56 AM, Ilmir Usmanov wrote:

> As usual, I'll mostly focus on fortran part, since I don't know CFE and
> ME well enough.
> 
> Nevertheless, is there a reason to differ C from Fortran in behavior of
> matching int-expr-list? I mean, in C you support only integer literals
> as parameter of wait, but in Fortran you also cover parameters.

Either Thomas or Jim Norris worked on the CFE. I just included in this
patch to make gomp-4_0-branch more in sync with our internal branch.

There are a couple of other places, like the parameter for
vector_length, where C only supports integer literals. I don't know what
the plan is for those.

> Anyway, Fortran part.
>>  gfc_match_oacc_wait (void)
>> +    if (!gfc_resolve_expr (el->expr)
>> +        || el->expr->ts.type != BT_INTEGER || el->expr->rank != 0
>> +        || el->expr->expr_type != EXPR_CONSTANT)
>> +      {
>> +        gfc_error ("WAIT clause at %L requires a scalar INTEGER
>> expression",
>> +               &el->expr->where);
>>
> As you probably know, in Fortran one can define a function after whole
> program.
> So, since you are resolving wait-list during match, you don't cover this
> case. But async clause works fine.

I moved that check back into resolve_omp_clauses. So it functions now.

> By the way, the same difference between wait and async presents in CFE.
>
> Otherwise, Fortran part looks good for me.

Thomas, any thoughts on the CFE?

Thanks,
Cesar


2014-10-06  Cesar Philippidis  <cesar@codesourcery.com>
	    James Norris  <jnorris@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	gcc/
	* builtin-types.def (BT_FN_VOID_INT_PTR_INT): Define.
	* oacc-builtins.def (DEF_GOACC_BUILTIN): Define.
	* omp-low.c (scan_sharing_clauses): Update handling of
	OMP_CLAUSE_ASYNC and OMP_CLAUSE_WAIT.
	(expand_oacc_offload): Likewise.
	(expand_omp_target): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses): Call gimplify_expr for
	OMP_CLAUSE_ASYNC and OMP_CLAUSE_WAIT.
	(gimplify_adjust_omp_clauses): Don't treat OMP_CLAUSE_ASYNC and
	OMP_CLAUSE_WAIT as unreachable.

	gcc/c-family/
	* c-common.h (c_finish_oacc_wait): Declare.
	* c-omp.c (c_finish_oacc_wait): New function.
	* c-pragma.c (oacc_pragmas): Add an entry for "wait".
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_WAIT.
	(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_ASYNC and
	PRAGMA_OMP_CLAUSE_WAIT.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Handle async and wait.
	(c_parser_oacc_integer_list): New function.
	(c_parser_oacc_int_list_parens): New function.
	(c_parser_oacc_clause_async): New function.
	(c_parser_oacc_clause_wait): New function.
	(c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_ASYNC and
	PRAGMA_OMP_CLAUSE_WAIT.
	(OACC_KERNELS_CLAUSE_MASK): Add async and wait clauses.
	(OACC_PARALLEL_CLAUSE_MASK): Likewise.
	(OACC_UPDATE_CLAUSE_MASK): Likewise.
	(OACC_WAIT_CLAUSE_MASK): New define.
	(c_parser_oacc_wait): New function.
	(c_parser_omp_construct): Handle PRAGMA_OACC_WAIT.
	* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_ASYNC and
	OMP_CLAUSE_WAIT.

	gcc/fortran/
	* gfortran.h (struct gfc_omp_clauses): Remove non_clause_wait_expr.
	* dump-parse-tree.c (show_omp_clauses): Likewise.
	* openmp.c (gfc_free_omp_clauses): Likewise.
	(gfc_match_omp_clauses): Update handling of async.
	(OACC_WAIT_CLAUSE_MASK): New define.
	(gfc_match_oacc_wait): Make the wait directive comply with OpenACC 2.0.
	(resolve_omp_clauses): Use resolve_oacc_scalar_in_expr inspect
	arguments to the wait clause.
	(resolve_oacc_wait): Remove.
	(gfc_resolve_oacc_directive): Handle EXEC_OACC_WAIT with
	resolve_omp_clauses.
	* trans-openmp.c (gfc_trans_omp_clauses): Update handling of OpenACC
	wait arguments.
	(gfc_trans_oacc_wait_directive): New function.
	(gfc_trans_oacc_directive): Use it.
	* types.def (BT_FN_VOID_INT_PTR_INT): Define.

	gcc/testsuite/
	* c-c++-common/goacc/asyncwait-1.c: New test.
	* gfortran.dg/goacc/asyncwait-1.f95: New test.
	* gfortran.dg/goacc/asyncwait-2.f95: New test.
	* gfortran.dg/goacc/asyncwait-3.f95: New test.
	* gfortran.dg/goacc/asyncwait-4.f95: New test.

	libgomp/
	* libgomp.map (GOACC_2.0): Add GOACC_wait.
	* libgomp_g.h (GOACC_wait): Declare.
	* oacc-parallel.c (GOACC_wait): Define.


diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 7c294af..094b3a8 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -358,6 +358,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_SIZE,
 		     BT_VOID, BT_PTR, BT_INT, BT_SIZE)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_INT,
 		     BT_VOID, BT_PTR, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_INT_PTR_INT,
+		     BT_VOID, BT_INT, BT_PTR, BT_INT)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_CONST_PTR_PTR_SIZE,
 		     BT_VOID, BT_CONST_PTR, BT_PTR, BT_SIZE)
 DEF_FUNCTION_TYPE_3 (BT_FN_INT_STRING_CONST_STRING_VALIST_ARG,
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 5ec79a0..a03b3ab 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1211,6 +1211,7 @@ extern void c_finish_omp_taskwait (location_t);
 extern void c_finish_omp_taskyield (location_t);
 extern tree c_finish_omp_for (location_t, enum tree_code, tree, tree, tree,
 			      tree, tree, tree);
+extern tree c_finish_oacc_wait (location_t, tree, tree);
 extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask,
 				 tree, tree *);
 extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree);
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 3c3fa44..ab417ad 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -29,7 +29,40 @@ along with GCC; see the file COPYING3.  If not see
 #include "c-pragma.h"
 #include "gimple-expr.h"
 #include "langhooks.h"
+#include "omp-low.h"
 
+/* Complete a #pragma oacc wait construct.  LOC is the location of
+   the #pragma.  */
+
+tree
+c_finish_oacc_wait (location_t loc, tree parms, tree clauses)
+{
+  const int nparms = list_length (parms);
+  tree stmt, t;
+  vec<tree, va_gc> *args;
+
+  vec_alloc (args, nparms + 2);
+  stmt = builtin_decl_explicit (BUILT_IN_GOACC_WAIT);
+
+  if (find_omp_clause (clauses, OMP_CLAUSE_ASYNC))
+    t = fold_convert (integer_type_node, OMP_CLAUSE_ASYNC_EXPR (clauses));
+  else
+    t = build_int_cst (integer_type_node, -2);  /* TODO: XXX FIX -2.  */
+
+  args->quick_push (t);
+  args->quick_push (build_int_cst (integer_type_node, nparms));
+
+  for (t = parms; t; t = TREE_CHAIN (t))
+    args->quick_push (build_int_cst (integer_type_node,
+			TREE_INT_CST_LOW (OMP_CLAUSE_WAIT_EXPR (t))));
+
+  stmt = build_call_expr_loc_vec (loc, stmt, args);
+  add_stmt (stmt);
+
+  vec_free (args);
+
+  return stmt;
+}
 
 /* Complete a #pragma omp master construct.  STMT is the structured-block
    that follows the pragma.  LOC is the l*/
diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index e3073bc..27e6c9f 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1177,6 +1177,7 @@ static const struct omp_pragma_def oacc_pragmas[] = {
   { "loop", PRAGMA_OACC_LOOP },
   { "parallel", PRAGMA_OACC_PARALLEL },
   { "update", PRAGMA_OACC_UPDATE },
+  { "wait", PRAGMA_OACC_WAIT },
 };
 static const struct omp_pragma_def omp_pragmas[] = {
   { "atomic", PRAGMA_OMP_ATOMIC },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index d83a700..ded496a 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -32,6 +32,7 @@ typedef enum pragma_kind {
   PRAGMA_OACC_LOOP,
   PRAGMA_OACC_PARALLEL,
   PRAGMA_OACC_UPDATE,
+  PRAGMA_OACC_WAIT,
   PRAGMA_OMP_ATOMIC,
   PRAGMA_OMP_BARRIER,
   PRAGMA_OMP_CANCEL,
@@ -76,6 +77,7 @@ typedef enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_NONE = 0,
 
   PRAGMA_OMP_CLAUSE_ALIGNED,
+  PRAGMA_OMP_CLAUSE_ASYNC,
   PRAGMA_OMP_CLAUSE_COLLAPSE,
   PRAGMA_OMP_CLAUSE_COPY,
   PRAGMA_OMP_CLAUSE_COPYIN,
@@ -127,6 +129,7 @@ typedef enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_UNIFORM,
   PRAGMA_OMP_CLAUSE_UNTIED,
   PRAGMA_OMP_CLAUSE_VECTOR_LENGTH,
+  PRAGMA_OMP_CLAUSE_WAIT,
   
   /* Clauses for Cilk Plus SIMD-enabled function.  */
   PRAGMA_CILK_CLAUSE_NOMASK,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index d118c37..534ff47 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -9750,6 +9750,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	case 'a':
 	  if (!strcmp ("aligned", p))
 	    result = PRAGMA_OMP_CLAUSE_ALIGNED;
+	  else if (!strcmp ("async", p))
+	    result = PRAGMA_OMP_CLAUSE_ASYNC;
 	  break;
 	case 'c':
 	  if (!strcmp ("collapse", p))
@@ -9887,6 +9889,10 @@ c_parser_omp_clause_name (c_parser *parser)
 	  else if (flag_cilkplus && !strcmp ("vectorlength", p))
 	    result = PRAGMA_CILK_CLAUSE_VECTORLENGTH;
 	  break;
+	case 'w':
+	  if (!strcmp ("wait", p))
+	    result = PRAGMA_OMP_CLAUSE_WAIT;
+	  break;
 	}
     }
 
@@ -9913,6 +9919,52 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
       }
 }
 
+/* OpenACC 2.0
+   integer-list:
+     integer
+     integer-list , integer
+
+   Parse a list of intergers.  */
+
+static tree
+c_parser_oacc_integer_list (c_parser *parser, location_t clause_loc,
+			    enum omp_clause_code kind, tree list)
+{
+  if (c_parser_peek_token (parser)->type == CPP_CLOSE_PAREN)
+    return list;
+
+  while (c_parser_peek_token (parser)->type == CPP_NUMBER)
+    {
+      tree t;
+
+      t = build_omp_clause (clause_loc, kind);
+
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (c_parser_peek_token (parser)->value)))
+	{
+	  c_parser_error (parser, "expected integer expression");
+	  return list;
+	}
+
+      OMP_CLAUSE_DECL (t) = c_parser_peek_token (parser)->value;
+      OMP_CLAUSE_CHAIN (t) = list;
+      list = t;
+      c_parser_consume_token (parser);
+
+      if (c_parser_peek_token (parser)->type == CPP_CLOSE_PAREN ||
+          !c_parser_require (parser, CPP_COMMA, "expected %<,%>"))
+	return list;
+
+      if (c_parser_peek_token (parser)->type != CPP_NUMBER)
+	{
+	  c_parser_error (parser, "expected integer expression");
+	  return list;
+	}
+    }
+
+    c_parser_error (parser, "expected integer expression");
+    return list;
+}
+
 /* OpenACC 2.0, OpenMP 2.5:
    variable-list:
      identifier
@@ -10019,6 +10071,21 @@ c_parser_omp_variable_list (c_parser *parser,
   return list;
 }
 
+static tree
+c_parser_oacc_int_list_parens (c_parser *parser, enum omp_clause_code kind,
+			      tree list)
+{
+  /* The clauses location.  */
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    {
+      list = c_parser_oacc_integer_list (parser, loc, kind, list);
+      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+    }
+  return list;
+}
+
 /* Similarly, but expect leading and trailing parenthesis.  This is a very
    common case for OpenACC and OpenMP clauses.  */
 
@@ -10497,6 +10564,96 @@ c_parser_omp_clause_num_workers (c_parser *parser, tree list)
   return list;
 }
 
+/* OpenACC:
+   async [( int-expr )] */
+
+static tree
+c_parser_oacc_clause_async (c_parser *parser, tree list)
+{
+  tree c, t = NULL_TREE;
+  location_t expr_loc, async_loc;
+
+  expr_loc = async_loc = c_parser_peek_token (parser)->location;
+  /* TODO XXX: FIX -1  (acc_async_noval).  */
+  t = build_int_cst (integer_type_node, -1);
+
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    {
+      bool error = false;
+      HOST_WIDE_INT n;
+
+      c_parser_consume_token (parser);
+      expr_loc = c_parser_peek_token (parser)->location;
+
+      if (c_parser_peek_token (parser)->type == CPP_NUMBER)
+	{
+	  t = c_parser_peek_token (parser)->value;
+	  t = c_fully_fold (t, false, NULL);
+
+	  if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
+	    || !tree_fits_shwi_p (t)
+	    || (n = tree_to_shwi (t)) <= -3		/* TODO XXX: FIX -3.  */
+	    || (int) n != n)
+	    {
+	      expr_loc = c_parser_peek_token (parser)->location;
+	      c_parser_error (parser, "expected integer expression");
+	      error = true;
+	    }
+	  else
+	    {
+	      c_parser_consume_token (parser);
+	    }
+	}
+      else
+	{
+	  c_parser_error (parser, "expected integer expression");
+	  error = true;
+	}
+
+      if (error ||
+		!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+	{
+	  return list;
+	}
+    }
+  else
+    {
+      t = c_fully_fold (t, false, NULL);
+    }
+
+  c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
+					build_int_cst (TREE_TYPE (t), 0));
+  if (CAN_HAVE_LOCATION_P (c))
+       SET_EXPR_LOCATION (c, expr_loc);
+  check_no_duplicate_clause (list, OMP_CLAUSE_ASYNC, "async");
+  c = build_omp_clause (async_loc, OMP_CLAUSE_ASYNC);
+  OMP_CLAUSE_ASYNC_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  list = c;
+
+  return list;
+}
+
+/* OpenACC:
+   wait ( int-expr-list ) */
+
+static tree
+c_parser_oacc_clause_wait (c_parser *parser, tree list)
+{
+  location_t clause_loc = c_parser_peek_token (parser)->location;
+
+  if (c_parser_peek_token (parser)->type != CPP_OPEN_PAREN)
+    return list;
+
+  c_parser_consume_token (parser);
+
+  list = c_parser_oacc_integer_list (parser, clause_loc, OMP_CLAUSE_WAIT, list);
+
+  c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+
+  return list;
+}
+
 /* OpenMP 2.5:
    ordered */
 
@@ -11354,6 +11511,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 
       switch (c_kind)
 	{
+	case PRAGMA_OMP_CLAUSE_ASYNC:
+	  clauses = c_parser_oacc_clause_async (parser, clauses);
+	  c_name = "async";
+	  break;
 	case PRAGMA_OMP_CLAUSE_COLLAPSE:
 	  clauses = c_parser_omp_clause_collapse (parser, clauses);
 	  c_name = "collapse";
@@ -11434,6 +11595,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_vector_length (parser, clauses);
 	  c_name = "vector_length";
 	  break;
+	case PRAGMA_OMP_CLAUSE_WAIT:
+	  clauses = c_parser_oacc_clause_wait (parser, clauses);
+	  c_name = "wait";
+	  break;
 	default:
 	  c_parser_error (parser, "expected clause");
 	  goto saw_error;
@@ -11748,7 +11913,8 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
 */
 
 #define OACC_KERNELS_CLAUSE_MASK					\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY)			\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE)		\
@@ -11758,7 +11924,8 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT)	\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
 
 static tree
 c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
@@ -11828,7 +11995,8 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
 */
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY)			\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE)		\
@@ -11842,7 +12010,8 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_VECTOR_LENGTH) )
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_VECTOR_LENGTH)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
 
 static tree
 c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
@@ -11881,10 +12050,12 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
 */
 
 #define OACC_UPDATE_CLAUSE_MASK						\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)		\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HOST)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)			\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF) )
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
 
 static void
 c_parser_oacc_update (c_parser *parser)
@@ -11910,6 +12081,30 @@ c_parser_oacc_update (c_parser *parser)
   add_stmt (stmt);
 }
 
+/* OpenACC 2.0:
+   # pragma acc wait [(intseq)] oacc-wait-clause[optseq] new-line
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_WAIT_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) )
+
+static tree
+c_parser_oacc_wait (location_t loc, c_parser *parser, char *p_name)
+{
+  tree stmt, clauses, list = NULL_TREE;
+
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    list = c_parser_oacc_int_list_parens (parser, OMP_CLAUSE_WAIT, list);
+
+  strcpy (p_name, " wait");
+  clauses = c_parser_oacc_all_clauses (parser, OACC_WAIT_CLAUSE_MASK, p_name);
+  stmt = c_finish_oacc_wait (loc, list, clauses);
+
+  return stmt;
+}
+
 /* OpenMP 2.5:
    # pragma omp atomic new-line
      expression-stmt
@@ -14248,6 +14443,10 @@ c_parser_omp_construct (c_parser *parser)
       strcpy (p_name, "#pragma acc");
       stmt = c_parser_oacc_parallel (loc, parser, p_name);
       break;
+    case PRAGMA_OACC_WAIT:
+      strcpy (p_name, "#pragma wait");
+      stmt = c_parser_oacc_wait (loc, parser, p_name);
+      break;
     case PRAGMA_OMP_ATOMIC:
       c_parser_omp_atomic (loc, parser);
       return;
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index ac036c3..7e95182 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12294,6 +12294,8 @@ c_finish_omp_clauses (tree clauses)
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  pc = &OMP_CLAUSE_CHAIN (c);
 	  continue;
 
diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index d7f2182..f85f6b6 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -1173,12 +1173,6 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
 	  fputc (')', dumpfile);
 	}
     }
-  if (omp_clauses->non_clause_wait_expr)
-    {
-      fputc ('(', dumpfile);
-      show_expr (omp_clauses->non_clause_wait_expr);
-      fputc (')', dumpfile);
-    }
   if (omp_clauses->sched_kind != OMP_SCHED_NONE)
     {
       const char *type;
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 63fb537..a8c2e81 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1264,7 +1264,6 @@ typedef struct gfc_omp_clauses
   struct gfc_expr *num_gangs_expr;
   struct gfc_expr *num_workers_expr;
   struct gfc_expr *vector_length_expr;
-  struct gfc_expr *non_clause_wait_expr;
   gfc_expr_list *wait_list;
   gfc_expr_list *tile_list;
   unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 4a48335..8e6f0ec 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -83,7 +83,6 @@ gfc_free_omp_clauses (gfc_omp_clauses *c)
   gfc_free_expr (c->num_gangs_expr);
   gfc_free_expr (c->num_workers_expr);
   gfc_free_expr (c->vector_length_expr);
-  gfc_free_expr (c->non_clause_wait_expr);
   for (i = 0; i < OMP_LIST_NUM; i++)
     gfc_free_omp_namelist (c->lists[i]);
   gfc_free_expr_list (c->wait_list);
@@ -496,10 +495,15 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask,
 	if (gfc_match ("async") == MATCH_YES)
 	  {
 	    c->async = true;
-	    if (gfc_match (" ( %e )", &c->async_expr) == MATCH_YES)
-	      needs_space = false;
-	    else
-	      needs_space = true;
+	    needs_space = false;
+	    if (gfc_match (" ( %e )", &c->async_expr) != MATCH_YES)
+	      {
+		c->async_expr = gfc_get_constant_expr (BT_INTEGER,
+						       gfc_default_integer_kind,
+						      &gfc_current_locus);
+		/* TODO XXX: FIX -1 (acc_async_noval).  */
+		mpz_set_si (c->async_expr->value.integer, -1);
+	      }
 	    continue;
 	  }
       if ((mask & OMP_CLAUSE_GANG) && !c->gang)
@@ -1168,6 +1172,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask,
 #define OACC_EXIT_DATA_CLAUSES \
   (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT | OMP_CLAUSE_COPYOUT \
    | OMP_CLAUSE_DELETE)
+#define OACC_WAIT_CLAUSES \
+  (OMP_CLAUSE_ASYNC)
 
 
 match
@@ -1328,8 +1334,31 @@ match
 gfc_match_oacc_wait (void)
 {
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
-  gfc_match (" ( %e )", &c->non_clause_wait_expr);
+  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);
+
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk in !$ACC WAIT at %C");
+      return MATCH_ERROR;
+    }
 
+  if (wait_list)
+    for (el = wait_list; el; el = el->next)
+      {
+	if (el->expr == NULL)
+	  {
+	    gfc_error ("Invalid argument to $!ACC WAIT at %L",
+		       &wait_list->expr->where);
+	    return MATCH_ERROR;
+	  }
+      }
+
+
+  c->wait = true;
+  c->wait_list = wait_list;
   new_st.op = EXEC_OACC_WAIT;
   new_st.ext.omp_clauses = c;
   return MATCH_YES;
@@ -3343,7 +3372,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
   if (omp_clauses->wait)
     if (omp_clauses->wait_list)
       for (el = omp_clauses->wait_list; el; el = el->next)
-	resolve_oacc_positive_int_expr (el->expr, "WAIT");
+	resolve_oacc_scalar_int_expr (el->expr, "WAIT");
 }
 
 
@@ -4490,16 +4519,6 @@ resolve_oacc_cache (gfc_code *code)
 }
 
 
-static void
-resolve_oacc_wait (gfc_code *code)
-{
-  gfc_expr_list* el;
-
-  for (el = code->ext.omp_clauses->wait_list; el; el = el->next)
-    resolve_oacc_positive_int_expr (el->expr, "WAIT");
-}
-
-
 void
 gfc_resolve_oacc_declare (gfc_namespace *ns)
 {
@@ -4573,6 +4592,7 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
     case EXEC_OACC_UPDATE:
     case EXEC_OACC_ENTER_DATA:
     case EXEC_OACC_EXIT_DATA:
+    case EXEC_OACC_WAIT:
       resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL,
 			   true);
       break;
@@ -4584,9 +4604,6 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
     case EXEC_OACC_CACHE:
       resolve_oacc_cache (code);
       break;
-    case EXEC_OACC_WAIT:
-      resolve_oacc_wait (code);
-      break;
     default:
       break;
     }
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index b32d857..87d1c94 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2545,6 +2545,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       c = build_omp_clause (where.lb->location, OMP_CLAUSE_INDEPENDENT);
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
+  if (clauses->wait_list)
+    {
+      gfc_expr_list *el;
+      tree list = NULL;
+
+      for (el = clauses->wait_list; el; el = el->next)
+	{
+	  c = build_omp_clause (where.lb->location, OMP_CLAUSE_WAIT);
+	  OMP_CLAUSE_DECL (c) = gfc_convert_expr_to_tree (block, el->expr);
+	  OMP_CLAUSE_CHAIN (c) = list;
+	  list = c;
+	}
+
+      omp_clauses = list;
+    }
   if (clauses->num_gangs_expr)
     {
       tree num_gangs_var = 
@@ -2617,14 +2632,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	  omp_clauses = gfc_trans_add_clause (c, omp_clauses);
 	}
     }
-  if (clauses->non_clause_wait_expr)
-    {
-      tree wait_var = 
-	  gfc_convert_expr_to_tree (block, clauses->non_clause_wait_expr);
-      c = build_omp_clause (where.lb->location, OMP_CLAUSE_WAIT);
-      OMP_CLAUSE_WAIT_EXPR (c)= wait_var;
-      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
-    }
 
   return nreverse (omp_clauses);
 }
@@ -2690,7 +2697,7 @@ gfc_trans_oacc_construct (gfc_code *code)
   return gfc_finish_block (&block);
 }
 
-/* update, enter_data, exit_data, wait, cache. */
+/* update, enter_data, exit_data, cache. */
 static tree 
 gfc_trans_oacc_executable_directive (gfc_code *code)
 {
@@ -2728,6 +2735,44 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
   return gfc_finish_block (&block);
 }
 
+static tree
+gfc_trans_oacc_wait_directive (gfc_code *code)
+{
+  stmtblock_t block;
+  tree stmt, t;
+  vec<tree, va_gc> *args;
+  int nparms = 0;
+  gfc_expr_list *el;
+  gfc_omp_clauses *clauses = code->ext.omp_clauses;
+  location_t loc = input_location;
+
+  for (el = clauses->wait_list; el; el = el->next)
+    nparms++;
+
+  vec_alloc (args, nparms + 2);
+  stmt = builtin_decl_explicit (BUILT_IN_GOACC_WAIT);
+
+  gfc_start_block (&block);
+
+  if (clauses->async_expr)
+    t = gfc_convert_expr_to_tree (&block, clauses->async_expr);
+  else
+    t = build_int_cst (integer_type_node, -2);
+
+  args->quick_push (t);
+  args->quick_push (build_int_cst (integer_type_node, nparms));
+
+  for (el = clauses->wait_list; el; el = el->next)
+    args->quick_push (gfc_convert_expr_to_tree (&block, el->expr));
+
+  stmt = build_call_expr_loc_vec (loc, stmt, args);
+  gfc_add_expr_to_block (&block, stmt);
+
+  vec_free (args);
+
+  return gfc_finish_block (&block);
+}
+
 static tree gfc_trans_omp_sections (gfc_code *, gfc_omp_clauses *);
 static tree gfc_trans_omp_workshare (gfc_code *, gfc_omp_clauses *);
 
@@ -4333,11 +4378,12 @@ gfc_trans_oacc_directive (gfc_code *code)
       return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
 			       NULL);
     case EXEC_OACC_UPDATE:
-    case EXEC_OACC_WAIT:
     case EXEC_OACC_CACHE:
     case EXEC_OACC_ENTER_DATA:
     case EXEC_OACC_EXIT_DATA:
       return gfc_trans_oacc_executable_directive (code);
+    case EXEC_OACC_WAIT:
+      return gfc_trans_oacc_wait_directive (code);
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 6c2fdc0..1dce308 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -145,6 +145,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I2_INT, BT_VOID, BT_VOLATILE_PTR, BT_I2, BT
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, BT_INT)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I8_INT, BT_VOID, BT_VOLATILE_PTR, BT_I8, BT_INT)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I16_INT, BT_VOID, BT_VOLATILE_PTR, BT_I16, BT_INT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_INT_PTR_INT, BT_VOID, BT_INT, BT_PTR, BT_INT)
 
 DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT,
                      BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 8e7bc26..5a8904f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6283,6 +6283,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
 			     is_gimple_val, fb_rvalue) == GS_ERROR)
 	    remove = true;
@@ -6293,8 +6295,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_GANG:
-	case OMP_CLAUSE_ASYNC:
-	case OMP_CLAUSE_WAIT:
 	case OMP_NO_CLAUSE_CACHE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_WORKER:
@@ -6690,6 +6690,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  break;
 
 	case OMP_CLAUSE_HOST:
@@ -6697,8 +6699,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_GANG:
-	case OMP_CLAUSE_ASYNC:
-	case OMP_CLAUSE_WAIT:
 	case OMP_NO_CLAUSE_CACHE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_WORKER:
diff --git a/gcc/oacc-builtins.def b/gcc/oacc-builtins.def
index e4bc756..1962a0f 100644
--- a/gcc/oacc-builtins.def
+++ b/gcc/oacc-builtins.def
@@ -39,5 +39,7 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
 		   ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
 		   BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait",
+		   BT_FN_VOID_INT_PTR_INT, ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
 			    BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index bb39f00..303c274 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1906,6 +1906,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_COLLAPSE:
 	  break;
 
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
+	  gcc_assert (is_gimple_omp_oacc_specifically (ctx->stmt));
+	  break;
+
 	case OMP_CLAUSE_ALIGNED:
 	  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
 	  decl = OMP_CLAUSE_DECL (c);
@@ -1919,8 +1924,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_GANG:
-	case OMP_CLAUSE_ASYNC:
-	case OMP_CLAUSE_WAIT:
 	case OMP_NO_CLAUSE_CACHE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_WORKER:
@@ -2055,11 +2058,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
 	  /* FALLTHRU */
+	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_COLLAPSE:
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
+	case OMP_CLAUSE_WAIT:
 	  break;
 
 	case OMP_CLAUSE_HOST:
@@ -2067,8 +2072,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_GANG:
-	case OMP_CLAUSE_ASYNC:
-	case OMP_CLAUSE_WAIT:
 	case OMP_NO_CLAUSE_CACHE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_WORKER:
@@ -5497,7 +5500,7 @@ expand_oacc_offload (struct omp_region *region)
 
   /* Emit a library call to launch CHILD_FN.  */
   tree t1, t2, t3, t4,
-    t_num_gangs, t_num_workers, t_vector_length,
+    t_num_gangs, t_num_workers, t_vector_length, t_async,
     device, cond, c, clauses;
   enum built_in_function start_ix;
   location_t clause_loc;
@@ -5522,6 +5525,8 @@ expand_oacc_offload (struct omp_region *region)
   t_num_gangs = t_num_workers = t_vector_length
     = fold_convert_loc (gimple_location (entry_stmt),
 			integer_type_node, integer_one_node);
+  t_async = fold_convert_loc (gimple_location (entry_stmt),
+			integer_type_node, build_int_cst (integer_type_node, -2));
   switch (region->type)
     {
     case GIMPLE_OACC_PARALLEL:
@@ -5542,6 +5547,13 @@ expand_oacc_offload (struct omp_region *region)
 	t_vector_length = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
 					    integer_type_node,
 					    OMP_CLAUSE_VECTOR_LENGTH_EXPR (c));
+      /* FALL THROUGH.  */
+    case GIMPLE_OACC_KERNELS:
+      c = find_omp_clause (clauses, OMP_CLAUSE_ASYNC);
+      if (c)
+	t_async = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+					    integer_type_node,
+					    OMP_CLAUSE_ASYNC_EXPR (c));
       break;
 
     default:
@@ -5643,10 +5655,48 @@ expand_oacc_offload (struct omp_region *region)
   gimple g;
   tree openmp_target = get_offload_symbol_decl ();
   tree fnaddr = build_fold_addr_expr (child_fn);
-  g = gimple_build_call (builtin_decl_explicit (start_ix), 10, device,
-			 fnaddr, build_fold_addr_expr (openmp_target),
-			 t1, t2, t3, t4,
-			 t_num_gangs, t_num_workers, t_vector_length);
+
+  vec<tree> *args;
+  int idx;
+
+  vec_alloc (args, 12);
+  args->quick_push (device);
+  args->quick_push (fnaddr);
+  args->quick_push (build_fold_addr_expr (openmp_target));
+  args->quick_push (t1);
+  args->quick_push (t2);
+  args->quick_push (t3);
+  args->quick_push (t4);
+  args->quick_push (t_num_gangs);
+  args->quick_push (t_num_workers);
+  args->quick_push (t_vector_length);
+  args->quick_push (t_async);
+  idx = args->length ();
+  args->quick_push (fold_convert_loc (gimple_location (entry_stmt),
+			integer_type_node, integer_minus_one_node));
+  c = find_omp_clause (clauses, OMP_CLAUSE_WAIT);
+  if (c)
+    {
+      int n = 0;
+
+      for (t = c; t; t = OMP_CLAUSE_CHAIN (t))
+	{
+	  if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_WAIT)
+	    {
+	      args->safe_push (fold_convert (integer_type_node,
+				OMP_CLAUSE_WAIT_EXPR (t)));
+	      n++;
+	    }
+	}
+
+        args->ordered_remove (idx);
+	args->quick_insert (idx, fold_convert_loc (gimple_location (entry_stmt),
+				 integer_type_node,
+				 build_int_cst (integer_type_node, n)));
+    }
+
+  g = gimple_build_call_vec (builtin_decl_explicit (start_ix), *args);
+  args->release ();
   gimple_set_location (g, gimple_location (entry_stmt));
   gsi_insert_before (&gsi, g, GSI_SAME_STMT);
 }
@@ -9379,17 +9429,63 @@ expand_omp_target (struct omp_region *region)
 
   gimple g;
   tree openmp_target = get_offload_symbol_decl ();
-  if (kind == GF_OMP_TARGET_KIND_REGION)
+  vec<tree> *args;
+
+  vec_alloc (args, 6);
+  args->quick_push (device);
+
+  if (kind ==  GF_OMP_TARGET_KIND_REGION)
+    args->quick_push (build_fold_addr_expr (child_fn));
+
+  args->quick_push (build_fold_addr_expr (openmp_target));
+  args->quick_push (t1);
+  args->quick_push (t2);
+  args->quick_push (t3);
+  args->safe_push (t4);
+
+  if (kind == GF_OMP_TARGET_KIND_OACC_DATA ||
+      kind == GF_OMP_TARGET_KIND_OACC_UPDATE)
     {
-      tree fnaddr = build_fold_addr_expr (child_fn);
-      g = gimple_build_call (builtin_decl_explicit (start_ix), 7, device,
-			     fnaddr, build_fold_addr_expr (openmp_target),
-			     t1, t2, t3, t4);
+      int idx;
+
+      c = find_omp_clause (clauses, OMP_CLAUSE_ASYNC);
+      if (c)
+	t1 = fold_convert_loc (OMP_CLAUSE_LOCATION (c), integer_type_node,
+				OMP_CLAUSE_ASYNC_EXPR (c));
+      else /* TODO: XXX FIX -2.  */
+	t1 = fold_convert_loc (gimple_location (entry_stmt),
+		      integer_type_node, build_int_cst (integer_type_node, -2));
+
+      args->safe_push (t1);
+      idx = args->length ();
+      args->safe_push (fold_convert_loc (gimple_location (entry_stmt),
+			integer_type_node, integer_minus_one_node));
+
+      c = find_omp_clause (clauses, OMP_CLAUSE_WAIT);
+      if (c)
+	{
+	  int n = 0;
+
+	  for (t = c; t; t = OMP_CLAUSE_CHAIN (t))
+	    {
+	      if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_WAIT)
+		{
+		  args->safe_push (fold_convert (integer_type_node,
+				OMP_CLAUSE_WAIT_EXPR (t)));
+		  n++;
+		}
+	    }
+
+	    args->ordered_remove (idx);
+	    args->quick_insert (idx,
+				fold_convert_loc (gimple_location (entry_stmt),
+				integer_type_node,
+				build_int_cst (integer_type_node, n)));
+	}
     }
-  else
-    g = gimple_build_call (builtin_decl_explicit (start_ix), 6, device,
-			   build_fold_addr_expr (openmp_target),
-			   t1, t2, t3, t4);
+
+  g = gimple_build_call_vec (builtin_decl_explicit (start_ix), *args);
+  args->release ();
   gimple_set_location (g, gimple_location (entry_stmt));
   gsi_insert_before (&gsi, g, GSI_SAME_STMT);
   if (kind != GF_OMP_TARGET_KIND_REGION)
diff --git a/gcc/testsuite/c-c++-common/goacc/asyncwait-1.c b/gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
new file mode 100644
index 0000000..0f7d297
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
@@ -0,0 +1,290 @@
+/* { dg-do compile } */
+
+void *malloc (__SIZE_TYPE__);
+
+int
+main (int argc, char **argv)
+{
+    int N = 64;
+    float *a, *b;
+    int i;
+
+    a = (float *) malloc (N * sizeof (float));
+    b = (float *) malloc (N * sizeof (float));
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,) /* { dg-error "expected '\\)' before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (,1) /* { dg-error "expected integer expression before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,) /* { dg-error "expected '\\)' before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2 3) /* { dg-error "expected '\\)' before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,,) /* { dg-error "expected '\\)' before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 /* { dg-error "expected '\\)' before end of line" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (*) /* { dg-error "expected integer expression before '\\*' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (a) /* { dg-error "expected integer expression before 'a'" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (N) /* { dg-error "expected integer expression before 'N'" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1.0) /* { dg-error "expected integer expression before numeric constant" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async () /* { dg-error "expected integer expression before '\\)' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 2) /* { dg-error "expected ',' before numeric constant" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,) /* { dg-error "expected integer expression before '\\)' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (,1) /* { dg-error "expected integer expression before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,) /* { dg-error "expected integer expression before '\\)' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2 3) /* { dg-error "expected ',' before numeric constant" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,,) /* { dg-error "expected integer expression before ',' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 /* { dg-error "expected ',' before end of line" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,*) /* { dg-error "expected integer expression before '\\*' token" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,a) /* { dg-error "expected integer expression before 'a'" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (a) /* { dg-error "expected integer expression before 'a'" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (N) /* { dg-error "expected integer expression before 'N'" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1.0) /* { dg-error "expected integer expression before numeric constant" } */
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait ()
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc wait (1 2) /* { dg-error "expected ',' before numeric constant" } */
+
+#pragma acc wait (1,) /* { dg-error "expected integer expression before '\\)' token" } */
+
+#pragma acc wait (,1) /* { dg-error "expected integer expression before ',' token" } */
+
+#pragma acc wait (1,2,) /* { dg-error "expected integer expression before '\\)' token" } */
+
+#pragma acc wait (1,2 3) /* { dg-error "expected ',' before numeric constant" } */
+
+#pragma acc wait (1,2,,) /* { dg-error "expected integer expression before ',' token" } */
+
+#pragma acc wait (1 /* { dg-error "expected ',' before end of line" } */
+
+#pragma acc wait (1,*) /* { dg-error "expected integer expression before '\\*' token" } */
+
+#pragma acc wait (1,a) /* { dg-error "expected integer expression before 'a'" } */
+
+#pragma acc wait (a) /* { dg-error "expected integer expression before 'a'" } */
+
+#pragma acc wait (N) /* { dg-error "expected integer expression before 'N'" } */
+
+#pragma acc wait (1.0) /* { dg-error "expected integer expression before numeric constant" } */
+
+#pragma acc wait 1 /* { dg-error "expected clause before numeric constant" } */
+
+#pragma acc wait N /* { dg-error "expected clause before 'N'" } */
+
+#pragma acc wait async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
+
+#pragma acc wait async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
+
+#pragma acc wait async (1,) /* { dg-error "expected '\\)' before ',' token" } */
+
+#pragma acc wait async (,1) /* { dg-error "expected integer expression before ',' token" } */
+
+#pragma acc wait async (1,2,) /* { dg-error "expected '\\)' before ',' token" } */
+
+#pragma acc wait async (1,2 3) /* { dg-error "expected '\\)' before ',' token" } */
+
+#pragma acc wait async (1,2,,) /* { dg-error "expected '\\)' before ',' token" } */
+
+#pragma acc wait async (1 /* { dg-error "expected '\\)' before end of line" } */
+
+#pragma acc wait async (*) /* { dg-error "expected integer expression before '\\*' token" } */
+
+#pragma acc wait async (a) /* { dg-error "expected integer expression before 'a'" } */
+
+#pragma acc wait async (N) /* { dg-error "expected integer expression before 'N'" } */
+
+#pragma acc wait async (1.0) /* { dg-error "expected integer expression before numeric constant" } */
+
+    return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-1.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-1.f95
new file mode 100644
index 0000000..d630d38
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-1.f95
@@ -0,0 +1,91 @@
+! { dg-do compile }
+
+program asyncwait
+  integer, parameter :: N = 64
+  real, allocatable :: a(:), b(:)
+  integer i
+
+  allocate (a(N))
+  allocate (b(N))
+
+  a(:) = 3.0
+  b(:) = 0.0
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) async (1 2) ! { 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)) async (1,) ! { 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)) async (,1) ! { dg-error "Invalid character in name" }
+  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)) async (1,2,) ! { 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)) async (1,2 3) ! { 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)) async (1,2,,) ! { 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)) async (1  ! { 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)) async (*) ! { dg-error "Invalid character in name at" }
+  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)) async (a) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) async (N)
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) async (1.0) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) async () ! { dg-error "Invalid character in name at " }
+  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)) async
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+end program asyncwait
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
new file mode 100644
index 0000000..db0ce1f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
@@ -0,0 +1,91 @@
+! { dg-do compile }
+
+program asyncwait
+  integer, parameter :: N = 64
+  real, allocatable :: a(:), b(:)
+  integer i
+
+  allocate (a(N))
+  allocate (b(N))
+
+  a(:) = 3.0
+  b(:) = 0.0
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait (1 2) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (1,) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (,1) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (1,2,) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (1,2 3) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (1,2,,) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (1 ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (*) ! { dg-error "Syntax error in OpenACC expression list" }
+  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 (a) ! { dg-error "WAIT clause at \\\(1\\\) requires a scalar INTEGER expression" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait (N)
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait (1.0) ! { dg-error "WAIT clause at \\\(1\\\) requires a scalar INTEGER expression" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait () ! { dg-error "Syntax error in OpenACC expression list" }
+  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)
+  end do
+  !$acc end parallel
+end program asyncwait
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
new file mode 100644
index 0000000..32c11de
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
@@ -0,0 +1,41 @@
+! { dg-do compile }
+
+program asyncwait
+  integer, parameter :: N = 64
+  real, allocatable :: a(:), b(:)
+  integer i
+
+  allocate (a(N))
+  allocate (b(N))
+
+  a(:) = 3.0
+  b(:) = 0.0
+
+  !$acc wait (1 2) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait (1,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait (,1) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait (1, 2, ) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait (1, 2, ,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait (1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait (1, *) ! { dg-error "Invalid argument to \\\$\\\!ACC WAIT" }
+
+  !$acc wait (1, a) ! { dg-error "WAIT clause at \\\(1\\\) requires a scalar INTEGER expression" }
+
+  !$acc wait (a) ! { dg-error "WAIT clause at \\\(1\\\) requires a scalar INTEGER expression" }
+
+  !$acc wait (N) 
+
+  !$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 N ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$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
new file mode 100644
index 0000000..cd64ef3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
@@ -0,0 +1,37 @@
+! { dg-do compile }
+
+program asyncwait
+  integer, parameter :: N = 64
+  real, allocatable :: a(:), b(:)
+  integer i
+
+  allocate (a(N))
+  allocate (b(N))
+
+  a(:) = 3.0
+  b(:) = 0.0
+
+  !$acc wait async (1 2) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (1,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (,1) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (1, 2, ) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (1, 2, ,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (1, *) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (1, a) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc wait async (a) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
+
+  !$acc wait async (N)
+
+  !$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" }
+end program asyncwait
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 69a4d83..382128d 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -246,4 +246,5 @@ GOACC_2.0 {
 	GOACC_kernels;
 	GOACC_parallel;
 	GOACC_update;
+	GOACC_wait;
 };
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 9dca76a9..f8a8d4b 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -225,5 +225,6 @@ extern void GOACC_kernels (int, void (*) (void *), const void *,
 extern void GOACC_parallel (int, void (*) (void *), const void *,
 			    size_t, void **, size_t *, unsigned short *,
 			    int, int, int);
+extern void GOACC_wait (int, int, ...);
 
 #endif /* LIBGOMP_G_H */
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 02fbb12..68ce728 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -138,3 +138,9 @@ acc_on_device (acc_device_t dev)
   return __builtin_acc_on_device (dev);
 }
 ialias (acc_on_device)
+
+void
+GOACC_wait (int async, int num_waits, ...)
+{
+  gomp_fatal ("Sorry, GOACC_wait is unimplemented.");
+}

Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]