This is the mail archive of the
mailing list for the GCC project.
[PATCH] OpenACC for C front end
- From: James Norris <jnorris at codesourcery dot com>
- To: <gcc-patches at gcc dot gnu dot org>
- Date: Wed, 5 Nov 2014 15:39:44 -0600
- Subject: [PATCH] OpenACC for C front end
- Authentication-results: sourceware.org; auth=none
Hi!
This patch represents the changes for OpenACC 2.0
in the C front-end. At present these files will
not compile as the changes for the middle end are
not present.
OK to commit?
Thanks,
Jim
=> c/ChangeLog
2014-11-05 James Norris <jnorris@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Ilmir Usmanov <i.usmanov@samsung.com>
* c-typeck.c (c_finish_oacc_parallel, c_finish_oacc_kernels,
c_finish_oacc_data): New functions.
(handle_omp_array_sections, c_finish_omp_clauses):
Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR, OMP_CLAUSE_NUM_GANGS,
OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_ASYNC,
and OMP_CLAUSE_WAIT.
(c_finish_omp_clauses): Handle PRAGMA_OMP_CLAUSE_NUM_GANGS,
PRAGMA_OMP_CLAUSE_NUM_WORKERS, and PRAGMA_OMP_CLAUSE_VECTOR_LENGTH.
* c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels,
c_finish_oacc_data): New prototypes.
* c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA,
PRAGMA_OACC_EXIT_DATA, and PRAGMA_OACC_UPDATE.
(c_parser_omp_clause_name): Handle OpenACC clauses.
(c_parser_oacc_wait_list, c_parser_oacc_data_clause,
c_parser_oacc_data_clause_deviceptr, c_parser_omp_clause_num_gangs,
c_parser_omp_clause_num_workers, c_parser_oacc_clause_async,
c_parser_oacc_clause_wait, c_parser_omp_clause_vector_length,
c_parser_oacc_all_clauses, c_parser_oacc_data, c_parser_oacc_kernels,
c_parser_oacc_enter_exit_data, c_parser_oacc_loop, c_parser_oacc_parallel,
c_parser_oacc_update, c_parser_oacc_wait): New functions.
(c_parser_omp_construct): Handle PRAGMA_OMP_DATA, PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL and PRAMGA_OACC_WAIT.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index d316216..df4b4cb 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1239,10 +1239,15 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
vec<tree, va_gc> **, location_t *,
tree *, vec<location_t> *,
unsigned int * = NULL);
+static tree c_parser_oacc_loop (location_t, c_parser *, char *);
static void c_parser_omp_construct (c_parser *);
static void c_parser_omp_threadprivate (c_parser *);
+static void c_parser_oacc_enter_exit_data (c_parser *, bool);
+static void c_parser_oacc_update (c_parser *);
static void c_parser_omp_barrier (c_parser *);
static void c_parser_omp_flush (c_parser *);
+static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code,
+ tree, tree *);
static void c_parser_omp_taskwait (c_parser *);
static void c_parser_omp_taskyield (c_parser *);
static void c_parser_omp_cancel (c_parser *);
@@ -4482,6 +4487,14 @@ c_parser_initval (c_parser *parser, struct c_expr *after,
Although they are erroneous if the labels declared aren't defined,
is it useful for the syntax to be this way?
+ OpenACC:
+
+ block-item:
+ openacc-directive
+
+ openacc-directive:
+ update-directive
+
OpenMP:
block-item:
@@ -4828,6 +4841,29 @@ c_parser_label (c_parser *parser)
@throw expression ;
@throw ;
+ OpenACC:
+
+ statement:
+ openacc-construct
+
+ openacc-construct:
+ parallel-construct
+ kernels-construct
+ data-construct
+ loop-construct
+
+ parallel-construct:
+ parallel-directive structured-block
+
+ kernels-construct:
+ kernels-directive structured-block
+
+ data-construct:
+ data-directive structured-block
+
+ loop-construct:
+ loop-directive structured-block
+
OpenMP:
statement:
@@ -9509,6 +9545,25 @@ c_parser_pragma (c_parser *parser, enum pragma_context context)
switch (id)
{
+ case PRAGMA_OACC_ENTER_DATA:
+ c_parser_oacc_enter_exit_data (parser, true);
+ return false;
+
+ case PRAGMA_OACC_EXIT_DATA:
+ c_parser_oacc_enter_exit_data (parser, false);
+ return false;
+
+ case PRAGMA_OACC_UPDATE:
+ if (context != pragma_compound)
+ {
+ if (context == pragma_stmt)
+ c_parser_error (parser, "%<#pragma acc update%> may only be "
+ "used in compound statements");
+ goto bad_stmt;
+ }
+ c_parser_oacc_update (parser);
+ return false;
+
case PRAGMA_OMP_BARRIER:
if (context != pragma_compound)
{
@@ -9711,7 +9766,7 @@ c_parser_pragma_pch_preprocess (c_parser *parser)
c_common_pch_pragma (parse_in, TREE_STRING_POINTER (name));
}
-/* OpenMP 2.5 / 3.0 / 3.1 / 4.0 parsing routines. */
+/* OpenACC and OpenMP parsing routines. */
/* Returns name of the next clause.
If the clause is not recognized PRAGMA_OMP_CLAUSE_NONE is returned and
@@ -9738,20 +9793,32 @@ 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))
result = PRAGMA_OMP_CLAUSE_COLLAPSE;
+ else if (!strcmp ("copy", p))
+ result = PRAGMA_OMP_CLAUSE_COPY;
else if (!strcmp ("copyin", p))
result = PRAGMA_OMP_CLAUSE_COPYIN;
+ else if (!strcmp ("copyout", p))
+ result = PRAGMA_OMP_CLAUSE_COPYOUT;
else if (!strcmp ("copyprivate", p))
result = PRAGMA_OMP_CLAUSE_COPYPRIVATE;
+ else if (!strcmp ("create", p))
+ result = PRAGMA_OMP_CLAUSE_CREATE;
break;
case 'd':
- if (!strcmp ("depend", p))
+ if (!strcmp ("delete", p))
+ result = PRAGMA_OMP_CLAUSE_DELETE;
+ else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
+ else if (!strcmp ("deviceptr", p))
+ result = PRAGMA_OMP_CLAUSE_DEVICEPTR;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -9763,6 +9830,10 @@ c_parser_omp_clause_name (c_parser *parser)
else if (!strcmp ("from", p))
result = PRAGMA_OMP_CLAUSE_FROM;
break;
+ case 'h':
+ if (!strcmp ("host", p))
+ result = PRAGMA_OMP_CLAUSE_SELF;
+ break;
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
@@ -9786,10 +9857,14 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
else if (!strcmp ("nowait", p))
result = PRAGMA_OMP_CLAUSE_NOWAIT;
+ else if (!strcmp ("num_gangs", p))
+ result = PRAGMA_OMP_CLAUSE_NUM_GANGS;
else if (!strcmp ("num_teams", p))
result = PRAGMA_OMP_CLAUSE_NUM_TEAMS;
else if (!strcmp ("num_threads", p))
result = PRAGMA_OMP_CLAUSE_NUM_THREADS;
+ else if (!strcmp ("num_workers", p))
+ result = PRAGMA_OMP_CLAUSE_NUM_WORKERS;
else if (flag_cilkplus && !strcmp ("nomask", p))
result = PRAGMA_CILK_CLAUSE_NOMASK;
break;
@@ -9800,6 +9875,20 @@ c_parser_omp_clause_name (c_parser *parser)
case 'p':
if (!strcmp ("parallel", p))
result = PRAGMA_OMP_CLAUSE_PARALLEL;
+ else if (!strcmp ("present", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT;
+ else if (!strcmp ("present_or_copy", p)
+ || !strcmp ("pcopy", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY;
+ else if (!strcmp ("present_or_copyin", p)
+ || !strcmp ("pcopyin", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN;
+ else if (!strcmp ("present_or_copyout", p)
+ || !strcmp ("pcopyout", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT;
+ else if (!strcmp ("present_or_create", p)
+ || !strcmp ("pcreate", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE;
else if (!strcmp ("private", p))
result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (!strcmp ("proc_bind", p))
@@ -9820,6 +9909,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_SHARED;
else if (!strcmp ("simdlen", p))
result = PRAGMA_OMP_CLAUSE_SIMDLEN;
+ else if (!strcmp ("self", p))
+ result = PRAGMA_OMP_CLAUSE_SELF;
break;
case 't':
if (!strcmp ("taskgroup", p))
@@ -9836,9 +9927,15 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
break;
case 'v':
- if (flag_cilkplus && !strcmp ("vectorlength", p))
+ if (!strcmp ("vector_length", p))
+ result = PRAGMA_OMP_CLAUSE_VECTOR_LENGTH;
+ 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;
}
}
@@ -9865,7 +9962,57 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
}
}
-/* OpenMP 2.5:
+/* OpenACC 2.0
+ Parse wait clause or wait directive parameters. */
+
+static tree
+c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
+{
+ vec<tree, va_gc> *args;
+ tree t, args_tree;
+
+ if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ return list;
+
+ args = c_parser_expr_list (parser, false, true, NULL, NULL, NULL, NULL);
+
+ if (args->length () == 0)
+ {
+ c_parser_error (parser, "expected integer expression before ')'");
+ release_tree_vector (args);
+ return list;
+ }
+
+ args_tree = build_tree_list_vec (args);
+
+ for (t = args_tree; t; t = TREE_CHAIN (t))
+ {
+ tree targ = TREE_VALUE (t);
+
+ if (targ != error_mark_node)
+ {
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (targ)))
+ {
+ c_parser_error (parser, "expression must be integral");
+ targ = error_mark_node;
+ }
+ else
+ {
+ tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
+
+ OMP_CLAUSE_DECL (c) = targ;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+ }
+ }
+ }
+
+ release_tree_vector (args);
+ c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+ return list;
+}
+
+/* OpenACC 2.0, OpenMP 2.5:
variable-list:
identifier
variable-list , identifier
@@ -9972,7 +10119,7 @@ c_parser_omp_variable_list (c_parser *parser,
}
/* Similarly, but expect leading and trailing parenthesis. This is a very
- common case for omp clauses. */
+ common case for OpenACC and OpenMP clauses. */
static tree
c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
@@ -9989,7 +10136,121 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
return list;
}
-/* OpenMP 3.0:
+/* OpenACC 2.0:
+ copy ( variable-list )
+ copyin ( variable-list )
+ copyout ( variable-list )
+ create ( variable-list )
+ delete ( variable-list )
+ present ( variable-list )
+ present_or_copy ( variable-list )
+ pcopy ( variable-list )
+ present_or_copyin ( variable-list )
+ pcopyin ( variable-list )
+ present_or_copyout ( variable-list )
+ pcopyout ( variable-list )
+ present_or_create ( variable-list )
+ pcreate ( variable-list ) */
+
+static tree
+c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
+ tree list)
+{
+ enum omp_clause_map_kind kind;
+ switch (c_kind)
+ {
+ default:
+ gcc_unreachable ();
+ case PRAGMA_OMP_CLAUSE_COPY:
+ kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ kind = OMP_CLAUSE_MAP_FORCE_ALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_DELETE:
+ kind = OMP_CLAUSE_MAP_FORCE_DEALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICE:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_HOST:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ kind = OMP_CLAUSE_MAP_FORCE_PRESENT;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ kind = OMP_CLAUSE_MAP_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ kind = OMP_CLAUSE_MAP_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
+ kind = OMP_CLAUSE_MAP_ALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_SELF:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ }
+ tree nl, c;
+ nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
+
+ for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_MAP_KIND (c) = kind;
+
+ return nl;
+}
+
+/* OpenACC 2.0:
+ deviceptr ( variable-list ) */
+
+static tree
+c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
+{
+ location_t loc = c_parser_peek_token (parser)->location;
+ tree vars, t;
+
+ /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic
+ c_parser_oacc_data_clause), as for PRAGMA_OMP_CLAUSE_DEVICEPTR,
+ variable-list must only allow for pointer variables. */
+ vars = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_ERROR, NULL);
+ for (t = vars; t && t; t = TREE_CHAIN (t))
+ {
+ tree v = TREE_PURPOSE (t);
+
+ /* FIXME diagnostics: Ideally we should keep individual
+ locations for all the variables in the var list to make the
+ following errors more precise. Perhaps
+ c_parser_omp_var_list_parens() should construct a list of
+ locations to go along with the var list. */
+
+ if (TREE_CODE (v) != VAR_DECL)
+ error_at (loc, "%qD is not a variable", v);
+ else if (TREE_TYPE (v) == error_mark_node)
+ ;
+ else if (!POINTER_TYPE_P (TREE_TYPE (v)))
+ error_at (loc, "%qD is not a pointer variable", v);
+
+ tree u = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (u) = OMP_CLAUSE_MAP_FORCE_DEVICEPTR;
+ OMP_CLAUSE_DECL (u) = v;
+ OMP_CLAUSE_CHAIN (u) = list;
+ list = u;
+ }
+
+ return list;
+}
+
+/* OpenACC 2.0, OpenMP 3.0:
collapse ( constant-expression ) */
static tree
@@ -10132,7 +10393,7 @@ c_parser_omp_clause_final (c_parser *parser, tree list)
return list;
}
-/* OpenMP 2.5:
+/* OpenACC, OpenMP 2.5:
if ( expression ) */
static tree
@@ -10200,6 +10461,51 @@ c_parser_omp_clause_nowait (c_parser *parser ATTRIBUTE_UNUSED, tree list)
return c;
}
+/* OpenACC:
+ num_gangs ( expression ) */
+
+static tree
+c_parser_omp_clause_num_gangs (c_parser *parser, tree list)
+{
+ location_t num_gangs_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ {
+ location_t expr_loc = c_parser_peek_token (parser)->location;
+ tree c, t = c_parser_expression (parser).value;
+ mark_exp_read (t);
+ t = c_fully_fold (t, false, NULL);
+
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ c_parser_error (parser, "expected integer expression");
+ return list;
+ }
+
+ /* Attempt to statically determine when the number isn't positive. */
+ 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);
+ if (c == boolean_true_node)
+ {
+ warning_at (expr_loc, 0,
+ "%<num_gangs%> value must be positive");
+ t = integer_one_node;
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NUM_GANGS, "num_gangs");
+
+ c = build_omp_clause (num_gangs_loc, OMP_CLAUSE_NUM_GANGS);
+ OMP_CLAUSE_NUM_GANGS_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+ }
+
+ return list;
+}
+
/* OpenMP 2.5:
num_threads ( expression ) */
@@ -10245,6 +10551,103 @@ c_parser_omp_clause_num_threads (c_parser *parser, tree list)
return list;
}
+/* OpenACC:
+ num_workers ( expression ) */
+
+static tree
+c_parser_omp_clause_num_workers (c_parser *parser, tree list)
+{
+ location_t num_workers_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ {
+ location_t expr_loc = c_parser_peek_token (parser)->location;
+ tree c, t = c_parser_expression (parser).value;
+ mark_exp_read (t);
+ t = c_fully_fold (t, false, NULL);
+
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ c_parser_error (parser, "expected integer expression");
+ return list;
+ }
+
+ /* Attempt to statically determine when the number isn't positive. */
+ 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);
+ if (c == boolean_true_node)
+ {
+ warning_at (expr_loc, 0,
+ "%<num_workers%> value must be positive");
+ t = integer_one_node;
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NUM_WORKERS, "num_workers");
+
+ c = build_omp_clause (num_workers_loc, OMP_CLAUSE_NUM_WORKERS);
+ OMP_CLAUSE_NUM_WORKERS_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+ }
+
+ return list;
+}
+
+/* OpenACC:
+ async [( int-expr )] */
+
+static tree
+c_parser_oacc_clause_async (c_parser *parser, tree list)
+{
+ tree c, t;
+ location_t 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)
+ {
+ c_parser_consume_token (parser);
+
+ t = c_parser_expression (parser).value;
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ c_parser_error (parser, "expected integer expression");
+ else if (t == error_mark_node
+ || !c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+ return list;
+ }
+ else
+ {
+ t = c_fully_fold (t, false, NULL);
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_ASYNC, "async");
+
+ c = build_omp_clause (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)
+ list = c_parser_oacc_wait_list (parser, clause_loc, list);
+
+ return list;
+}
+
/* OpenMP 2.5:
ordered */
@@ -10496,33 +10899,78 @@ c_parser_omp_clause_untied (c_parser *parser ATTRIBUTE_UNUSED, tree list)
return c;
}
-/* OpenMP 4.0:
- inbranch
- notinbranch */
+/* OpenACC:
+ vector_length ( expression ) */
static tree
-c_parser_omp_clause_branch (c_parser *parser ATTRIBUTE_UNUSED,
- enum omp_clause_code code, tree list)
+c_parser_omp_clause_vector_length (c_parser *parser, tree list)
{
- check_no_duplicate_clause (list, code, omp_clause_code_name[code]);
-
- tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
- OMP_CLAUSE_CHAIN (c) = list;
+ location_t vector_length_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ {
+ location_t expr_loc = c_parser_peek_token (parser)->location;
+ tree c, t = c_parser_expression (parser).value;
+ mark_exp_read (t);
+ t = c_fully_fold (t, false, NULL);
- return c;
-}
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
-/* OpenMP 4.0:
- parallel
- for
- sections
- taskgroup */
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ c_parser_error (parser, "expected integer expression");
+ return list;
+ }
-static tree
-c_parser_omp_clause_cancelkind (c_parser *parser ATTRIBUTE_UNUSED,
- enum omp_clause_code code, tree list)
-{
- tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
+ /* Attempt to statically determine when the number isn't positive. */
+ 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);
+ if (c == boolean_true_node)
+ {
+ warning_at (expr_loc, 0,
+ "%<vector_length%> value must be positive");
+ t = integer_one_node;
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH, "vector_length");
+
+ c = build_omp_clause (vector_length_loc, OMP_CLAUSE_VECTOR_LENGTH);
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+ }
+
+ return list;
+}
+
+/* OpenMP 4.0:
+ inbranch
+ notinbranch */
+
+static tree
+c_parser_omp_clause_branch (c_parser *parser ATTRIBUTE_UNUSED,
+ enum omp_clause_code code, tree list)
+{
+ check_no_duplicate_clause (list, code, omp_clause_code_name[code]);
+
+ tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+}
+
+/* OpenMP 4.0:
+ parallel
+ for
+ sections
+ taskgroup */
+
+static tree
+c_parser_omp_clause_cancelkind (c_parser *parser ATTRIBUTE_UNUSED,
+ enum omp_clause_code code, tree list)
+{
+ tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
OMP_CLAUSE_CHAIN (c) = list;
return c;
@@ -11032,6 +11480,144 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list)
return list;
}
+/* Parse all OpenACC clauses. The set clauses allowed by the directive
+ is a bitmask in MASK. Return the list of clauses found. */
+
+static tree
+c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
+ const char *where, bool finish_p = true)
+{
+ tree clauses = NULL;
+ bool first = true;
+
+ while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+ {
+ location_t here;
+ pragma_omp_clause c_kind;
+ const char *c_name;
+ tree prev = clauses;
+
+ if (!first && c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+
+ here = c_parser_peek_token (parser)->location;
+ c_kind = c_parser_omp_clause_name (parser);
+
+ 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";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPY:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYOUT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "create";
+ break;
+ case PRAGMA_OMP_CLAUSE_DELETE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "delete";
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "device";
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICEPTR:
+ clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
+ c_name = "deviceptr";
+ break;
+ case PRAGMA_OMP_CLAUSE_HOST:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "host";
+ break;
+ case PRAGMA_OMP_CLAUSE_IF:
+ clauses = c_parser_omp_clause_if (parser, clauses);
+ c_name = "if";
+ break;
+ case PRAGMA_OMP_CLAUSE_NUM_GANGS:
+ clauses = c_parser_omp_clause_num_gangs (parser, clauses);
+ c_name = "num_gangs";
+ break;
+ case PRAGMA_OMP_CLAUSE_NUM_WORKERS:
+ clauses = c_parser_omp_clause_num_workers (parser, clauses);
+ c_name = "num_workers";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_create";
+ break;
+ case PRAGMA_OMP_CLAUSE_REDUCTION:
+ clauses = c_parser_omp_clause_reduction (parser, clauses);
+ c_name = "reduction";
+ break;
+ case PRAGMA_OMP_CLAUSE_SELF:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "self";
+ break;
+ case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH:
+ 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;
+ }
+
+ first = false;
+
+ if (((mask >> c_kind) & 1) == 0 && !parser->error)
+ {
+ /* Remove the invalid clause(s) from the list to avoid
+ confusing the rest of the compiler. */
+ clauses = prev;
+ error_at (here, "%qs is not valid for %qs", c_name, where);
+ }
+ }
+
+ saw_error:
+ c_parser_skip_to_pragma_eol (parser);
+
+ if (finish_p)
+ return c_finish_omp_clauses (clauses);
+
+ return clauses;
+}
+
/* Parse all OpenMP clauses. The set clauses allowed by the directive
is a bitmask in MASK. Return the list of clauses found; the result
of clause default goes in *pdefault. */
@@ -11262,7 +11848,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
return clauses;
}
-/* OpenMP 2.5:
+/* OpenACC 2.0, OpenMP 2.5:
structured-block:
statement
@@ -11278,6 +11864,326 @@ c_parser_omp_structured_block (c_parser *parser)
return pop_stmt_list (stmt);
}
+/* OpenACC 2.0:
+ # pragma acc data oacc-data-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_DATA_CLAUSE_MASK \
+ ( (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) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \
+ | (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) )
+
+static tree
+c_parser_oacc_data (location_t loc, c_parser *parser)
+{
+ tree stmt, clauses, block;
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+ "#pragma acc data");
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+
+ stmt = c_finish_oacc_data (loc, clauses, block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc kernels oacc-kernels-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_KERNELS_CLAUSE_MASK \
+ ( (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) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \
+ | (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_WAIT) )
+
+static tree
+c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
+{
+ tree stmt, clauses = NULL_TREE, block;
+
+ strcat (p_name, " kernels");
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if (strcmp (p, "loop") == 0)
+ {
+ c_parser_consume_token (parser);
+ block = c_begin_omp_parallel ();
+ c_parser_oacc_loop (loc, parser, p_name);
+ stmt = c_finish_oacc_kernels (loc, clauses, block);
+ OACC_KERNELS_COMBINED (stmt) = 1;
+ return stmt;
+ }
+ }
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
+ p_name);
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+
+ stmt = c_finish_oacc_kernels (loc, clauses, block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc enter data oacc-enter-data-clause[optseq] new-line
+
+ or
+
+ # pragma acc exit data oacc-exit-data-clause[optseq] new-line
+
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_ENTER_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
+
+#define OACC_EXIT_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
+
+static void
+c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
+{
+ location_t loc = c_parser_peek_token (parser)->location;
+ tree clauses, stmt;
+
+ c_parser_consume_pragma (parser);
+
+ if (!c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_parser_error (parser, enter
+ ? "expected %<data%> in %<#pragma acc enter data%>"
+ : "expected %<data%> in %<#pragma acc exit data%>");
+ c_parser_skip_to_pragma_eol (parser);
+ return;
+ }
+
+ const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if (strcmp (p, "data") != 0)
+ {
+ c_parser_error (parser, "invalid pragma");
+ c_parser_skip_to_pragma_eol (parser);
+ return;
+ }
+
+ c_parser_consume_token (parser);
+
+ if (enter)
+ clauses = c_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK,
+ "#pragma acc enter data");
+ else
+ clauses = c_parser_oacc_all_clauses (parser, OACC_EXIT_DATA_CLAUSE_MASK,
+ "#pragma acc exit data");
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (loc, enter
+ ? "%<#pragma acc enter data%> has no data movement clause"
+ : "%<#pragma acc exit data%> has no data movement clause");
+ return;
+ }
+
+ stmt = enter ? make_node (OACC_ENTER_DATA) : make_node (OACC_EXIT_DATA);;
+ TREE_TYPE (stmt) = void_type_node;
+ if (enter)
+ OACC_ENTER_DATA_CLAUSES (stmt) = clauses;
+ else
+ OACC_EXIT_DATA_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+ add_stmt (stmt);
+}
+
+
+/* OpenACC 2.0:
+
+ # pragma acc loop oacc-loop-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_LOOP_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) )
+
+static tree
+c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
+{
+ tree stmt, clauses, block;
+
+ strcat (p_name, " loop");
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
+
+ block = c_begin_compound_stmt (true);
+ stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
+ block = c_end_compound_stmt (loc, block, true);
+ add_stmt (block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc parallel oacc-parallel-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_PARALLEL_CLAUSE_MASK \
+ ( (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) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_GANGS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_WORKERS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \
+ | (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_REDUCTION) \
+ | (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)
+{
+ tree stmt, clauses = NULL_TREE, block;
+
+ strcat (p_name, " parallel");
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if (strcmp (p, "loop") == 0)
+ {
+ c_parser_consume_token (parser);
+ block = c_begin_omp_parallel ();
+ c_parser_oacc_loop (loc, parser, p_name);
+ stmt = c_finish_oacc_parallel (loc, clauses, block);
+ OACC_PARALLEL_COMBINED (stmt) = 1;
+ return stmt;
+ }
+ }
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
+ p_name);
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+
+ stmt = c_finish_oacc_parallel (loc, clauses, block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc update oacc-update-clause[optseq] new-line
+*/
+
+#define OACC_UPDATE_CLAUSE_MASK \
+ ( (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_WAIT) )
+
+static void
+c_parser_oacc_update (c_parser *parser)
+{
+ location_t loc = c_parser_peek_token (parser)->location;
+
+ c_parser_consume_pragma (parser);
+
+ tree clauses = c_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK,
+ "#pragma acc update");
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (loc,
+ "%<#pragma acc update%> must contain at least one "
+ "%<device%> or %<host/self%> clause");
+ return;
+ }
+
+ if (parser->error)
+ return;
+
+ tree stmt = make_node (OACC_UPDATE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_UPDATE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+ 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 clauses, list = NULL_TREE, stmt = NULL_TREE;
+
+ if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+ list = c_parser_oacc_wait_list (parser, loc, 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
@@ -11754,10 +12660,11 @@ c_parser_omp_flush (c_parser *parser)
c_finish_omp_flush (loc);
}
-/* Parse the restricted form of the for statement allowed by OpenMP.
+/* Parse the restricted form of loop statements allowed by OpenACC and OpenMP.
The real trick here is to determine the loop control variable early
so that we can push a new decl if necessary to make it private.
- LOC is the location of the OMP in "#pragma omp". */
+ LOC is the location of the "acc" or "omp" in "#pragma acc" or "#pragma omp",
+ respectively. */
static tree
c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
@@ -12010,6 +12917,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
if (cclauses != NULL
&& cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL] != NULL)
{
+ gcc_assert (code != OACC_LOOP);
tree *c;
for (c = &cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL]; *c ; )
if (OMP_CLAUSE_CODE (*c) != OMP_CLAUSE_FIRSTPRIVATE
@@ -13599,6 +14507,25 @@ c_parser_omp_construct (c_parser *parser)
switch (p_kind)
{
+ case PRAGMA_OACC_DATA:
+ stmt = c_parser_oacc_data (loc, parser);
+ break;
+ case PRAGMA_OACC_KERNELS:
+ strcpy (p_name, "#pragma acc");
+ stmt = c_parser_oacc_kernels (loc, parser, p_name);
+ break;
+ case PRAGMA_OACC_LOOP:
+ strcpy (p_name, "#pragma acc");
+ stmt = c_parser_oacc_loop (loc, parser, p_name);
+ break;
+ case PRAGMA_OACC_PARALLEL:
+ 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-tree.h b/gcc/c/c-tree.h
index f7e723b..bcfec28 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -640,6 +640,9 @@ extern tree c_finish_bc_stmt (location_t, tree *, bool);
extern tree c_finish_goto_label (location_t, tree);
extern tree c_finish_goto_ptr (location_t, tree);
extern tree c_expr_to_decl (tree, bool *, bool *);
+extern tree c_finish_oacc_parallel (location_t, tree, tree);
+extern tree c_finish_oacc_kernels (location_t, tree, tree);
+extern tree c_finish_oacc_data (location_t, tree, tree);
extern tree c_begin_omp_parallel (void);
extern tree c_finish_omp_parallel (location_t, tree, tree);
extern tree c_begin_omp_task (void);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 79dbc3d..76503e4 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -11230,6 +11230,63 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se)
return expr;
}
+/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_PARALLEL. */
+
+tree
+c_finish_oacc_parallel (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_PARALLEL);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_PARALLEL_CLAUSES (stmt) = clauses;
+ OACC_PARALLEL_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
+/* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_KERNELS. */
+
+tree
+c_finish_oacc_kernels (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_KERNELS);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_KERNELS_CLAUSES (stmt) = clauses;
+ OACC_KERNELS_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
+/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_DATA. */
+
+tree
+c_finish_oacc_data (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DATA_CLAUSES (stmt) = clauses;
+ OACC_DATA_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
/* Like c_begin_compound_stmt, except force the retention of the BLOCK. */
tree
@@ -11761,6 +11818,7 @@ handle_omp_array_sections (tree c)
OMP_CLAUSE_SIZE (c) = size;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
if (!c_mark_addressable (t))
@@ -11824,7 +11882,7 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data)
return NULL_TREE;
}
-/* For all elements of CLAUSES, validate them vs OpenMP constraints.
+/* For all elements of CLAUSES, validate them against their constraints.
Remove any elements from the list that are invalid. */
tree
@@ -12184,7 +12242,9 @@ c_finish_omp_clauses (tree clauses)
else if (!c_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)))
&& !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
@@ -12253,6 +12313,11 @@ c_finish_omp_clauses (tree clauses)
case OMP_CLAUSE_TASKGROUP:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ 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;
