On 10/03/2014 07:34 AM, Cesar Philippidis wrote:
> On 09/24/2014 12:18 AM, Ilmir Usmanov wrote:
>> Hi Cesar!
>>
>> Thank you for the patch!
>>
>> On 24.09.2014 02:29, Cesar Philippidis wrote:
>>> This patch adds support for the async clause in the wait directive in
>>> fortran. It should be pretty straight forward. The fortran FE already
>>> supports the wait directive, but the async clause was introduced to the
>>> wait directive in OpenACC 2.0 and that was missing in gomp-4_0-branch.
>> Yes, I've mostly focused on spec. ver. 1.0.
>>
>>> Is this OK for gomp-4_0-branch?
>> No, it isn't. According to the spec and this presentation:
>> http://www.pgroup.com/lit/presentations/cea-3.pdf (See slide 1-35)
>> it is possible to write construction like:
>> !$acc wait(1) async(2)
>> However, your patch doesn't support this. Also, don't forget to check
>> whether a queue waits itself (for example, wait(1) async(1)).
>> In addition, it breaks current support of the directive (for example,
>> wait(1)).
>
> Sorry for the delay. I encountered some problems with the runtime in our
> internal branch, and that slowed things down a bit.
>
> Anyway, you are correct, I broke the optional argument to wait in the
> previous patch. This new patch addresses that and it also make the wait
> construct conform with OpenACC 2.0. Specifically,
>
> !$acc wait (1, 2) async (3)
>
> should behave as ex[ected.
>
> If you look at gfc_trans_oacc_wait_directive, you'll note that a call to
> GOACC_wait is emitted for the wait directive. Since I had to add a
> runtime library stub for that builtin function, I decided to go ahead an
> include the c front end bits.
>
> Is this patch OK for gomp-4_0-branch? Julian is working on working on a
> more complete implementation of the runtime. The runtime stub that I
> included is only temporary.
I noticed that I forgot to include the changes to gimplify.c in that
patch. This new patch includes those changes.
Cesar
2014-10-04 Cesar Philippidis <ce...@codesourcery.com>
James Norris <jnor...@codesourcery.com>
Thomas Schwinge <tho...@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..c158128 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,38 @@ 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;
+ }
+
+ 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);
+ return MATCH_ERROR;
+ }
+ }
+ c->wait_list = wait_list;
new_st.op = EXEC_OACC_WAIT;
new_st.ext.omp_clauses = c;
return MATCH_YES;
@@ -3343,7 +3379,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 +4526,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 +4599,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 +4611,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.");
+}
