Hi!
On 11/13/2014 07:02 AM, Jakub Jelinek wrote:
On Wed, Nov 05, 2014 at 03:37:08PM -0600, James Norris wrote:
2014-11-05 James Norris <jnor...@codesourcery.com>
Cesar Philippidis <ce...@codesourcery.com>
Thomas Schwinge <tho...@codesourcery.com>
Ilmir Usmanov <i.usma...@samsung.com>
...
Please check formatting. I see various spots with 8 spaces instead of tabs,
e.g. on
+ check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH,
+ "vector_length", location);
even the alignment is wrong, I see calls without space before (:
+ if (args == NULL || args->length() == 0)
...
+ error("%<wait%> expression must be integral");
other spots where the alignment isn't right:
Fixed.
+static tree
+cp_parser_oacc_cache (cp_parser *parser,
+ cp_token *pragma_tok __attribute__((unused)))
(cp_token should be below cp_parser). While at this,
__attribute__((unused)) should be replaced by ATTRIBUTE_UNUSED, or removing
the parameter name, or removing the parameter altogether even better.
Fixed by removing unused parameter.
For the formatting issues, you can easily look for them in the patch
(in lines starting with +), change the patch and apply interdiff to your
tree.
Thank you for the pointer to interdiff!
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -29,8 +29,47 @@ 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"
As Thomas? has said, you should include gomp-constants.h and use them in:
+ t = build_int_cst (integer_type_node, -2); /* TODO: XXX FIX -2. */
spots like this.
Fixed.
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1180,6 +1180,16 @@ typedef struct
static vec<pragma_ns_name> registered_pp_pragmas;
struct omp_pragma_def { const char *name; unsigned int id; };
+static const struct omp_pragma_def oacc_pragmas[] = {
+ { "data", PRAGMA_OACC_DATA },
+ { "enter", PRAGMA_OACC_ENTER_DATA },
+ { "exit", PRAGMA_OACC_EXIT_DATA },
+ { "kernels", PRAGMA_OACC_KERNELS },
+ { "loop", PRAGMA_OACC_LOOP },
+ { "parallel", PRAGMA_OACC_PARALLEL },
+ { "update", PRAGMA_OACC_UPDATE },
+ { "wait", PRAGMA_OACC_WAIT },
I'd avoid the , after the last element.
Fixed.
@@ -1383,6 +1402,17 @@ c_invoke_pragma_handler (unsigned int id)
void
init_pragma (void)
{
+ if (flag_openacc)
+ {
+ const int n_oacc_pragmas
+ = sizeof (oacc_pragmas) / sizeof (*oacc_pragmas);
+ int i;
+
+ for (i = 0; i < n_oacc_pragmas; ++i)
+ cpp_register_deferred_pragma (parse_in, "acc", oacc_pragmas[i].name,
+ oacc_pragmas[i].id, true, true);
+ }
+
if (flag_openmp)
{
const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
Is -fopenmp -fopenacc tested not to run out of number of supported pragmas
by libcpp?
Tests will be added in a follow-up patch once the middle end has been
accepted.
@@ -65,23 +74,30 @@ typedef enum pragma_kind {
} pragma_kind;
-/* All clauses defined by OpenMP 2.5, 3.0, 3.1 and 4.0.
+/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, and 4.0.
Used internally by both C and C++ parsers. */
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,
+ PRAGMA_OMP_CLAUSE_COPYOUT,
PRAGMA_OMP_CLAUSE_COPYPRIVATE,
+ PRAGMA_OMP_CLAUSE_CREATE,
PRAGMA_OMP_CLAUSE_DEFAULT,
+ PRAGMA_OMP_CLAUSE_DELETE,
PRAGMA_OMP_CLAUSE_DEPEND,
PRAGMA_OMP_CLAUSE_DEVICE,
+ PRAGMA_OMP_CLAUSE_DEVICEPTR,
PRAGMA_OMP_CLAUSE_DIST_SCHEDULE,
PRAGMA_OMP_CLAUSE_FINAL,
PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_OMP_CLAUSE_FOR,
PRAGMA_OMP_CLAUSE_FROM,
+ PRAGMA_OMP_CLAUSE_HOST,
PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OMP_CLAUSE_INBRANCH,
PRAGMA_OMP_CLAUSE_LASTPRIVATE,
@@ -90,16 +106,24 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_MERGEABLE,
PRAGMA_OMP_CLAUSE_NOTINBRANCH,
PRAGMA_OMP_CLAUSE_NOWAIT,
+ PRAGMA_OMP_CLAUSE_NUM_GANGS,
PRAGMA_OMP_CLAUSE_NUM_TEAMS,
PRAGMA_OMP_CLAUSE_NUM_THREADS,
+ PRAGMA_OMP_CLAUSE_NUM_WORKERS,
PRAGMA_OMP_CLAUSE_ORDERED,
PRAGMA_OMP_CLAUSE_PARALLEL,
+ PRAGMA_OMP_CLAUSE_PRESENT,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE,
PRAGMA_OMP_CLAUSE_PRIVATE,
PRAGMA_OMP_CLAUSE_PROC_BIND,
PRAGMA_OMP_CLAUSE_REDUCTION,
PRAGMA_OMP_CLAUSE_SAFELEN,
PRAGMA_OMP_CLAUSE_SCHEDULE,
PRAGMA_OMP_CLAUSE_SECTIONS,
+ PRAGMA_OMP_CLAUSE_SELF,
PRAGMA_OMP_CLAUSE_SHARED,
PRAGMA_OMP_CLAUSE_SIMDLEN,
PRAGMA_OMP_CLAUSE_TASKGROUP,
@@ -107,6 +131,8 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_TO,
PRAGMA_OMP_CLAUSE_UNIFORM,
PRAGMA_OMP_CLAUSE_UNTIED,
+ PRAGMA_OMP_CLAUSE_VECTOR_LENGTH,
+ PRAGMA_OMP_CLAUSE_WAIT,
Like for CILK, I'd strongly prefer if for the clauses that are
specific to OpenACC only you'd use PRAGMA_OACC_CLAUSE_* instead,
and put them after the PRAGMA_CILK_* enum values.
If you want to have PRAGMA_OACC_CLAUSE_ aliases also for the
clauses shared in between OpenMP and OpenACC, feel free to add
aliases like Cilk+ has them.
Fixed by renaming OpenACC-specific clauses. Have deferred
on creating alias names.
It is unfortunately lots of new clauses
and we are getting close to the 64 clauses limit :( when they are used
in bitmasks.
Will put my 'thinking cap' on.
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -27437,6 +27437,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_IF;
else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DEFAULT))
result = PRAGMA_OMP_CLAUSE_DEFAULT;
+ else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DELETE))
+ result = PRAGMA_OMP_CLAUSE_DELETE;
else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_PRIVATE))
result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR))
@@ -27451,20 +27453,30 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'a':
if (!strcmp ("aligned", p))
result = PRAGMA_OMP_CLAUSE_ALIGNED;
+ else if (!strcmp ("async", p))
+ result = PRAGMA_OMP_CLAUSE_ASYNC;
Please adjust to PRAGMA_OACC_CLAUSE_* here etc. too; that will make it
more clear that the clauses are OpenACC specific.
Fix as described above.
+ /* 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. */
Please avoid writing todo lists into the source, that should be
tracked in some PR instead. It is a general thing, we really should
have some expression that would be used in the FEs
to have a FE only expression code that would wrap decls and constants
that don't have a location to provide the location for them, then
we can use it everywhere, including the omp variable lists.
We need it for better diagnostics (to have right locus or locus ranges
in FE diagnostics).
Will submit a PR.
+ check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH,
+ "vector_length", location);
See above about formatting.
Fixed.
+static tree
+cp_parser_oacc_wait_list (cp_parser *parser, location_t clause_loc, tree list)
+{
+ vec<tree, va_gc> *args;
+ tree t, args_tree;
+
+ args = cp_parser_parenthesized_expression_list (parser, non_attr,
+ /*cast_p=*/false,
+ /*allow_expansion_p=*/true,
+ /*non_constant_p=*/NULL);
+
+ if (args == NULL || args->length() == 0)
+ {
+ cp_parser_error (parser, "expected integer expression before ')'");
Is cp_parser_parenthesized_expression_list an integer expression? Sounds
the wording doesn't match the implementation. Also, please use
%<)%> instead of ')'.
Fixed.
+ tree targ = TREE_VALUE (t);
+
+ if (targ != error_mark_node)
+ {
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (targ)))
+ {
+ error("%<wait%> expression must be integral");
+ targ = error_mark_node;
Why do you assign to targ at all, when it is not used afterwards?
If you remove it, drop the {}s around error too. And see above for
formatting comment.
Fixed.
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
+ || !tree_fits_shwi_p (t)
+ || (n = tree_to_shwi (t)) <= 0
+ || (int) n != n)
+ {
+ error_at (location, "expected positive integer expression");
+ return list;
+ }
Do you use it as int value anywhere, or only pass through as tree to some
function? If the latter, then neither the tree_fits* nor (int) n != n check
is needed, all you should care about is INTEGRAL_TYPE_P and tree_int_cst_sgn
0. Also, does it really have to be INTEGER_CST already during parsing, or
is it enough if it is INTEGER_CST after folding/template instantiation etc.?
The reason why omp_clause_collapse does the above is that the integer is
directly used during the parsing. For other integers, it is usually checked
in finish_omp_clauses (after instantiation) or so.
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
+ || !tree_fits_shwi_p (t)
+ || (n = tree_to_shwi (t)) <= 0
+ || (int) n != n)
+ {
+ error_at (location, "expected positive integer expression");
+ return list;
+ }
Likewise.
Fixed both occurences.
+static tree
+cp_parser_oacc_clause_async (cp_parser *parser, tree list)
+{
+ tree c, t;
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+
+ /* TODO XXX: FIX -1 (acc_async_noval). */
Again, please use gomp-constants.h constant here.
+ t = build_int_cst (integer_type_node, -1);
Fixed.
+/* 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.
Don't understand the comment here, LOC isn't mentioned anywhere in there,
nor it is a parameter of the function.
Comment removed along with similar occurences.
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL)
+ || cp_lexer_next_token_is_not (parser->lexer, CPP_NAME))
+ {
+ cp_parser_error (parser, enter
+ ? "expected %<data%> in %<#pragma acc enter data%>"
+ : "expected %<data%> in %<#pragma acc exit data%>");
Shouldn't the diagnostics for
#pragma acc enter
#pragma acc enter for
#pragma acc enter data2
be all the same? The cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL)
test doesn't make any sense together with
|| cp_lexer_next_token_is_not (parser->lexer, CPP_NAME), for CPP_PRAGMA_EOL
the next token is not CPP_NAME.
Also, as #pragma acc enter data is the name of the directive, it shouldn't
be translated, so IMHO you can just use
"expected %<data%> after %<#pragma acc %s%>", enter ? "enter" : "exit".
Fixed.
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return NULL_TREE;
+ }
+
+ const char *p =
+ IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
+ if (strcmp (p, "data") != 0)
+ {
+ cp_parser_error (parser, "invalid pragma");
As written above, best set p to "", if
cp_lexer_next_token_is (parser->lexer, CPP_NAME) is true, set it to
IDENTIFIER_POINTER ... and then strcmp and use one diagnostics for that.
Fixed.
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (pragma_tok->location,
+ "%<#pragma acc enter data%> has no data movement clause");
Unconditional enter even when it could be exit?
Fixed.
+ case OMP_CLAUSE_ASYNC:
Wonder if similar thing to PRAGMA_OACC_CLAUSE_* shouldn't be used here,
either OACC_CLAUSE_ASYNC or OMP_CLAUSE_ACC_ASYNC to make it clear
it is OpenACC thing. Though in this case I don't feel as strong about it as
in case of PRAGMA_OACC_CLAUSE_*.
I don't have a strong feeling one way or another either, so I'll
leave the OMP_CLAUSES_* alone for the time being.
+ t = OMP_CLAUSE_ASYNC_EXPR (c);
+ if (t == error_mark_node)
+ remove = true;
+ else if (!type_dependent_expression_p (t)
+ && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ error ("%<async%> expression must be integral");
You have OMP_CLAUSE_LOCATION (c) which you could use for error_at.
I followed the convention that was used elsewhere in the function
at using error ().
Otherwise LGTM.
Thank you for taking the time to review!
OK to commit after middle end has been accepted?
Regards,
Jim
=> cp/ChangeLog
2014-XX-XX James Norris <jnor...@codesourcery.com>
Cesar Philippidis <ce...@codesourcery.com>
Thomas Schwinge <tho...@codesourcery.com>
Ilmir Usmanov <i.usma...@samsung.com>
* cp-tree.h (finish_oacc_data, finish_oacc_kernels,
finish_oacc_parallel): New prototypes.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_ASYNC,
OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_WAIT.
(finish_oacc_data, finish_oacc_kernels, finish_oacc_parallel): New
functions.
* parser.c (cp_parser_omp_clause_name): Don't parse the
identifier for RID_DELETE. Add parsing of OpenACC clauses.
(cp_parser_omp_var_list_no_open): Add handling of array specifier.
(cp_parser_oacc_data_clause, cp_parser_oacc_data_clause_deviceptr,
cp_parser_oacc_vector_length, cp_parser_oacc_wait_list,
cp_parser_oacc_clause_wait, cp_parser_omp_clause_num_gangs,
cp_parser_omp_clause_num_workers, cp_parser_oacc_clause_async,
cp_parser_oacc_all_clauses, cp_parser_oacc_cache, cp_parser_oacc_data,
cp_parser_oacc_enter_exit_data, cp_parser_oacc_kernels,
cp_parser_oacc_loop, cp_parser_oacc_parallel, cp_parser_oacc_update,
cp_parser_oacc_wait): New functions.
(cp_parser_omp_construct): Handle OpenACC pragmas.
(cp_parser_pragma): Likewise.
=> c-family/ChangeLog
2014-XX-XX James Norris <jnor...@codesourcery.com>
Cesar Philippidis <ce...@codesourcery.com>
Thomas Schwinge <tho...@codesourcery.com>
Ilmir Usmanov <i.usma...@samsung.com>
* c-pragma.h (pragma_kind): Add OpenACC pragmas.
(pragma_omp_clause): Add OpenACC clauses.
* c-pragma.c (oacc_pragmas): New pragma definition array.
(c_pp_lookup_pragma, init_pragma): Handle OpenACC pragmas.
* c-cppbuiltin.c (c_cpp_builtins): Conditionally define _OPENACC.
* c.opt (fopenacc): New option.
* c-omp.c (c_finish_oacc_wait): New function.
(c_omp_split_clauses): Catch OACC_PARALLEL.
* c-common.h (c_finish_oacc_wait): New prototype.
* c-common.c (DEF_FUNCTION_TYPE_8, DEF_FUNTION_TYPE_12): New defines.
diff --git a/gcc/c-family/c-common.c b/gcc/c-family/c-common.c
index 95b6b1b..a2b7360 100644
--- a/gcc/c-family/c-common.c
+++ b/gcc/c-family/c-common.c
@@ -5213,6 +5213,11 @@ enum c_builtin_type
#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
NAME,
+#define DEF_FUNCTION_TYPE_VAR_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) NAME,
+#define DEF_FUNCTION_TYPE_VAR_12(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, \
+ ARG12) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
#include "builtin-types.def"
#undef DEF_PRIMITIVE_TYPE
@@ -5231,6 +5236,8 @@ enum c_builtin_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
BT_LAST
};
@@ -5323,6 +5330,14 @@ c_define_builtins (tree va_list_ref_type_node, tree va_list_arg_type_node)
def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4);
#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5);
+#define DEF_FUNCTION_TYPE_VAR_8(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) \
+ def_fn_type (ENUM, RETURN, 1, 8, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8);
+#define DEF_FUNCTION_TYPE_VAR_12(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) \
+ def_fn_type (ENUM, RETURN, 1, 12, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8, ARG9, ARG10, ARG11, ARG12);
#define DEF_POINTER_TYPE(ENUM, TYPE) \
builtin_types[(int) ENUM] = build_pointer_type (builtin_types[(int) TYPE]);
@@ -5344,6 +5359,8 @@ c_define_builtins (tree va_list_ref_type_node, tree va_list_arg_type_node)
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 7e53923..2c3d3aa 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1232,6 +1232,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-cppbuiltin.c b/gcc/c-family/c-cppbuiltin.c
index c571d1b..47e464e 100644
--- a/gcc/c-family/c-cppbuiltin.c
+++ b/gcc/c-family/c-cppbuiltin.c
@@ -1208,6 +1208,9 @@ c_cpp_builtins (cpp_reader *pfile)
else if (flag_stack_protect == 1)
cpp_define (pfile, "__SSP__=1");
+ if (flag_openacc)
+ cpp_define (pfile, "_OPENACC=201306");
+
if (flag_openmp)
cpp_define (pfile, "_OPENMP=201307");
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index d6ca3df..e22cb4b 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -29,8 +29,47 @@ 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 = 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))
+ {
+ if (TREE_CODE (OMP_CLAUSE_WAIT_EXPR (t)) == INTEGER_CST)
+ args->quick_push (build_int_cst (integer_type_node,
+ TREE_INT_CST_LOW (OMP_CLAUSE_WAIT_EXPR (t))));
+ else
+ args->quick_push (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*/
@@ -664,6 +703,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
enum c_omp_clause_split s;
int i;
+ gcc_assert (code != OACC_PARALLEL);
for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
cclauses[i] = NULL;
/* Add implicit nowait clause on
diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index 3183410..6a0dbb3 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1180,6 +1180,17 @@ typedef struct
static vec<pragma_ns_name> registered_pp_pragmas;
struct omp_pragma_def { const char *name; unsigned int id; };
+static const struct omp_pragma_def oacc_pragmas[] = {
+ { "cache", PRAGMA_OACC_CACHE },
+ { "data", PRAGMA_OACC_DATA },
+ { "enter", PRAGMA_OACC_ENTER_DATA },
+ { "exit", PRAGMA_OACC_EXIT_DATA },
+ { "kernels", PRAGMA_OACC_KERNELS },
+ { "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 },
{ "barrier", PRAGMA_OMP_BARRIER },
@@ -1212,11 +1223,20 @@ static const struct omp_pragma_def omp_pragmas_simd[] = {
void
c_pp_lookup_pragma (unsigned int id, const char **space, const char **name)
{
+ const int n_oacc_pragmas = sizeof (oacc_pragmas) / sizeof (*oacc_pragmas);
const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
const int n_omp_pragmas_simd = sizeof (omp_pragmas_simd)
/ sizeof (*omp_pragmas);
int i;
+ for (i = 0; i < n_oacc_pragmas; ++i)
+ if (oacc_pragmas[i].id == id)
+ {
+ *space = "acc";
+ *name = oacc_pragmas[i].name;
+ return;
+ }
+
for (i = 0; i < n_omp_pragmas; ++i)
if (omp_pragmas[i].id == id)
{
@@ -1383,6 +1403,17 @@ c_invoke_pragma_handler (unsigned int id)
void
init_pragma (void)
{
+ if (flag_openacc)
+ {
+ const int n_oacc_pragmas
+ = sizeof (oacc_pragmas) / sizeof (*oacc_pragmas);
+ int i;
+
+ for (i = 0; i < n_oacc_pragmas; ++i)
+ cpp_register_deferred_pragma (parse_in, "acc", oacc_pragmas[i].name,
+ oacc_pragmas[i].id, true, true);
+ }
+
if (flag_openmp)
{
const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index b9f09ba..fd00e14 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -27,6 +27,15 @@ along with GCC; see the file COPYING3. If not see
typedef enum pragma_kind {
PRAGMA_NONE = 0,
+ PRAGMA_OACC_CACHE,
+ PRAGMA_OACC_DATA,
+ PRAGMA_OACC_ENTER_DATA,
+ PRAGMA_OACC_EXIT_DATA,
+ PRAGMA_OACC_KERNELS,
+ PRAGMA_OACC_LOOP,
+ PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_UPDATE,
+ PRAGMA_OACC_WAIT,
PRAGMA_OMP_ATOMIC,
PRAGMA_OMP_BARRIER,
PRAGMA_OMP_CANCEL,
@@ -65,7 +74,7 @@ typedef enum pragma_kind {
} pragma_kind;
-/* All clauses defined by OpenMP 2.5, 3.0, 3.1 and 4.0.
+/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, and 4.0.
Used internally by both C and C++ parsers. */
typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_NONE = 0,
@@ -118,7 +127,25 @@ typedef enum pragma_omp_clause {
PRAGMA_CILK_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_CILK_CLAUSE_LASTPRIVATE = PRAGMA_OMP_CLAUSE_LASTPRIVATE,
PRAGMA_CILK_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION,
- PRAGMA_CILK_CLAUSE_UNIFORM = PRAGMA_OMP_CLAUSE_UNIFORM
+ PRAGMA_CILK_CLAUSE_UNIFORM = PRAGMA_OMP_CLAUSE_UNIFORM,
+
+ /* Clauses specific to OpenACC. */
+ PRAGMA_OACC_CLAUSE_ASYNC,
+ PRAGMA_OACC_CLAUSE_COPY,
+ PRAGMA_OACC_CLAUSE_COPYOUT,
+ PRAGMA_OACC_CLAUSE_CREATE,
+ PRAGMA_OACC_CLAUSE_DELETE,
+ PRAGMA_OACC_CLAUSE_DEVICEPTR,
+ PRAGMA_OACC_CLAUSE_NUM_GANGS,
+ PRAGMA_OACC_CLAUSE_NUM_WORKERS,
+ PRAGMA_OACC_CLAUSE_PRESENT,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
+ PRAGMA_OACC_CLAUSE_SELF,
+ PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
+ PRAGMA_OACC_CLAUSE_WAIT
} pragma_omp_clause;
extern struct cpp_reader* parse_in;
diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt
index 85dcb98..ad6e237 100644
--- a/gcc/c-family/c.opt
+++ b/gcc/c-family/c.opt
@@ -1282,6 +1282,10 @@ fobjc-std=objc1
ObjC ObjC++ Var(flag_objc1_only)
Conform to the Objective-C 1.0 language as implemented in GCC 4.0
+fopenacc
+C ObjC C++ ObjC++ Var(flag_openacc)
+Enable OpenACC
+
fopenmp
C ObjC C++ ObjC++ Var(flag_openmp)
Enable OpenMP (implies -frecursive in Fortran)
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index b3781ab..53be646 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5969,6 +5969,9 @@ extern tree finish_omp_clauses (tree);
extern void finish_omp_threadprivate (tree);
extern tree begin_omp_structured_block (void);
extern tree finish_omp_structured_block (tree);
+extern tree finish_oacc_data (tree, tree);
+extern tree finish_oacc_kernels (tree, tree);
+extern tree finish_oacc_parallel (tree, tree);
extern tree begin_omp_parallel (void);
extern tree finish_omp_parallel (tree, tree);
extern tree begin_omp_task (void);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index b106f3b..726d3ab 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -55,6 +55,7 @@ along with GCC; see the file COPYING3. If not see
#include "parser.h"
#include "type-utils.h"
#include "omp-low.h"
+#include "gomp-constants.h"
/* The lexer. */
@@ -27431,6 +27432,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_IF;
else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DEFAULT))
result = PRAGMA_OMP_CLAUSE_DEFAULT;
+ else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DELETE))
+ result = PRAGMA_OACC_CLAUSE_DELETE;
else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_PRIVATE))
result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR))
@@ -27445,20 +27448,30 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'a':
if (!strcmp ("aligned", p))
result = PRAGMA_OMP_CLAUSE_ALIGNED;
+ else if (!strcmp ("async", p))
+ result = PRAGMA_OACC_CLAUSE_ASYNC;
break;
case 'c':
if (!strcmp ("collapse", p))
result = PRAGMA_OMP_CLAUSE_COLLAPSE;
+ else if (!strcmp ("copy", p))
+ result = PRAGMA_OACC_CLAUSE_COPY;
else if (!strcmp ("copyin", p))
result = PRAGMA_OMP_CLAUSE_COPYIN;
+ else if (!strcmp ("copyout", p))
+ result = PRAGMA_OACC_CLAUSE_COPYOUT;
else if (!strcmp ("copyprivate", p))
result = PRAGMA_OMP_CLAUSE_COPYPRIVATE;
+ else if (!strcmp ("create", p))
+ result = PRAGMA_OACC_CLAUSE_CREATE;
break;
case 'd':
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_OACC_CLAUSE_DEVICEPTR;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -27470,6 +27483,10 @@ cp_parser_omp_clause_name (cp_parser *parser)
else if (!strcmp ("from", p))
result = PRAGMA_OMP_CLAUSE_FROM;
break;
+ case 'h':
+ if (!strcmp ("host", p))
+ result = PRAGMA_OACC_CLAUSE_SELF;
+ break;
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
@@ -27495,10 +27512,14 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_NOWAIT;
else if (flag_cilkplus && !strcmp ("nomask", p))
result = PRAGMA_CILK_CLAUSE_NOMASK;
+ else if (!strcmp ("num_gangs", p))
+ result = PRAGMA_OACC_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_OACC_CLAUSE_NUM_WORKERS;
break;
case 'o':
if (!strcmp ("ordered", p))
@@ -27507,6 +27528,22 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'p':
if (!strcmp ("parallel", p))
result = PRAGMA_OMP_CLAUSE_PARALLEL;
+ else if (!strcmp ("present", p))
+ result = PRAGMA_OACC_CLAUSE_PRESENT;
+ else if (!strcmp ("present_or_copy", p)
+ || !strcmp ("pcopy", p))
+ result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
+ else if (!strcmp ("present_or_copyin", p)
+ || !strcmp ("pcopyin", p))
+ result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
+ else if (!strcmp ("present_or_copyout", p)
+ || !strcmp ("pcopyout", p))
+ result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
+ else if (!strcmp ("present_or_create", p)
+ || !strcmp ("pcreate", p))
+ result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
+ else if (!strcmp ("private", p))
+ result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (!strcmp ("proc_bind", p))
result = PRAGMA_OMP_CLAUSE_PROC_BIND;
break;
@@ -27521,6 +27558,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_SCHEDULE;
else if (!strcmp ("sections", p))
result = PRAGMA_OMP_CLAUSE_SECTIONS;
+ else if (!strcmp ("self", p))
+ result = PRAGMA_OACC_CLAUSE_SELF;
else if (!strcmp ("shared", p))
result = PRAGMA_OMP_CLAUSE_SHARED;
else if (!strcmp ("simdlen", p))
@@ -27541,9 +27580,15 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
break;
case 'v':
- if (flag_cilkplus && !strcmp ("vectorlength", p))
+ if (!strcmp ("vector_length", p))
+ result = PRAGMA_OACC_CLAUSE_VECTOR_LENGTH;
+ else if (flag_cilkplus && !strcmp ("vectorlength", p))
result = PRAGMA_CILK_CLAUSE_VECTORLENGTH;
break;
+ case 'w':
+ if (!strcmp ("wait", p))
+ result = PRAGMA_OACC_CLAUSE_WAIT;
+ break;
}
}
@@ -27619,6 +27664,14 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
{
switch (kind)
{
+ case OMP_NO_CLAUSE_CACHE:
+ if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_SQUARE)
+ {
+ error_at (token->location, "expected %<[%>");
+ decl = error_mark_node;
+ break;
+ }
+ /* FALL THROUGH. */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
@@ -27649,6 +27702,29 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
if (!cp_parser_require (parser, CPP_CLOSE_SQUARE,
RT_CLOSE_SQUARE))
goto skip_comma;
+
+ if (kind == OMP_NO_CLAUSE_CACHE)
+ {
+ mark_exp_read (low_bound);
+ mark_exp_read (length);
+
+ if (TREE_CODE (low_bound) != INTEGER_CST
+ && !TREE_READONLY (low_bound))
+ {
+ error_at (token->location,
+ "%qD is not a constant", low_bound);
+ decl = error_mark_node;
+ }
+
+ if (TREE_CODE (length) != INTEGER_CST
+ && !TREE_READONLY (length))
+ {
+ error_at (token->location,
+ "%qD is not a constant", length);
+ decl = error_mark_node;
+ }
+ }
+
decl = tree_cons (low_bound, length, decl);
}
break;
@@ -27711,6 +27787,217 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
return list;
}
+/* 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
+cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
+ tree list)
+{
+ enum omp_clause_map_kind kind;
+ switch (c_kind)
+ {
+ default:
+ gcc_unreachable ();
+ case PRAGMA_OACC_CLAUSE_COPY:
+ kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OACC_CLAUSE_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_CREATE:
+ kind = OMP_CLAUSE_MAP_FORCE_ALLOC;
+ break;
+ case PRAGMA_OACC_CLAUSE_DELETE:
+ kind = OMP_CLAUSE_MAP_FORCE_DEALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICE:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OACC_CLAUSE_SELF:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT:
+ kind = OMP_CLAUSE_MAP_FORCE_PRESENT;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
+ kind = OMP_CLAUSE_MAP_TOFROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
+ kind = OMP_CLAUSE_MAP_TO;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
+ kind = OMP_CLAUSE_MAP_ALLOC;
+ break;
+ }
+ tree nl, c;
+ nl = cp_parser_omp_var_list (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
+cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
+{
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree vars, t;
+
+ /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic
+ cp_parser_oacc_data_clause), as for PRAGMA_OMP_CLAUSE_DEVICEPTR,
+ variable-list must only allow for pointer variables. */
+ vars = cp_parser_omp_var_list (parser, OMP_CLAUSE_ERROR, NULL);
+ for (t = vars; t; t = TREE_CHAIN (t))
+ {
+ tree v = TREE_PURPOSE (t);
+
+ 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:
+ vector_length ( expression ) */
+
+static tree
+cp_parser_oacc_clause_vector_length (cp_parser *parser, tree list)
+{
+ tree t, c;
+ location_t location = cp_lexer_peek_token (parser->lexer)->location;
+ bool error = false;
+ HOST_WIDE_INT n;
+
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ return list;
+
+ t = cp_parser_condition (parser);
+ if (t == error_mark_node || !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ error_at (location, "expected positive integer expression");
+ error = true;
+ }
+
+ if (error || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH, "vector_length",
+ location);
+
+ c = build_omp_clause (location, OMP_CLAUSE_VECTOR_LENGTH);
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ return list;
+}
+
+/* OpenACC 2.0
+ Parse wait clause or directive parameters. */
+
+static tree
+cp_parser_oacc_wait_list (cp_parser *parser, location_t clause_loc, tree list)
+{
+ vec<tree, va_gc> *args;
+ tree t, args_tree;
+
+ args = cp_parser_parenthesized_expression_list (parser, non_attr,
+ /*cast_p=*/false,
+ /*allow_expansion_p=*/true,
+ /*non_constant_p=*/NULL);
+
+ if (args == NULL || args->length () == 0)
+ {
+ cp_parser_error (parser,
+ "expected integer expression list before '%<)%>'");
+ if (args != NULL)
+ release_tree_vector (args);
+ return list;
+ }
+
+ args_tree = build_tree_list_vec (args);
+
+ release_tree_vector (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)))
+ error ("%<wait%> expression must be integral");
+ else
+ {
+ tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
+
+ mark_rvalue_use (targ);
+ OMP_CLAUSE_DECL (c) = targ;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+ }
+ }
+ }
+
+ return list;
+}
+
+/* OpenACC:
+ wait ( int-expr-list ) */
+
+static tree
+cp_parser_oacc_clause_wait (cp_parser *parser, tree list)
+{
+ location_t location = cp_lexer_peek_token (parser->lexer)->location;
+
+ if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_PAREN)
+ return list;
+
+ list = cp_parser_oacc_wait_list (parser, location, list);
+
+ return list;
+}
+
/* OpenMP 3.0:
collapse ( constant-expression ) */
@@ -27899,6 +28186,43 @@ cp_parser_omp_clause_nowait (cp_parser * /*parser*/,
return c;
}
+/* OpenACC:
+ num_gangs ( expression ) */
+
+static tree
+cp_parser_omp_clause_num_gangs (cp_parser *parser, tree list)
+{
+ tree t, c;
+ location_t location = cp_lexer_peek_token (parser->lexer)->location;
+ HOST_WIDE_INT n;
+
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ return list;
+
+ t = cp_parser_condition (parser);
+
+ if (t == error_mark_node
+ || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ error_at (location, "expected positive integer expression");
+ return list;
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NUM_GANGS, "num_gangs", location);
+
+ c = build_omp_clause (location, 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 ) */
@@ -27929,6 +28253,47 @@ cp_parser_omp_clause_num_threads (cp_parser *parser, tree list,
return c;
}
+/* OpenACC:
+ num_workers ( expression ) */
+
+static tree
+cp_parser_omp_clause_num_workers (cp_parser *parser, tree list)
+{
+ tree t, c;
+ location_t location = cp_lexer_peek_token (parser->lexer)->location;
+ HOST_WIDE_INT n;
+
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ return list;
+
+ t = cp_parser_condition (parser);
+
+ if (t == error_mark_node
+ || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
+ || !tree_fits_shwi_p (t)
+ || (n = tree_to_shwi (t)) <= 0
+ || (int) n != n)
+ {
+ error_at (location, "expected positive integer expression");
+ return list;
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NUM_WORKERS, "num_gangs",
+ location);
+
+ c = build_omp_clause (location, OMP_CLAUSE_NUM_WORKERS);
+ OMP_CLAUSE_NUM_WORKERS_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ return list;
+}
+
/* OpenMP 2.5:
ordered */
@@ -28623,6 +28988,175 @@ cp_parser_omp_clause_proc_bind (cp_parser *parser, tree list,
return list;
}
+/* OpenACC:
+ async [( int-expr )] */
+
+static tree
+cp_parser_oacc_clause_async (cp_parser *parser, tree list)
+{
+ tree c, t;
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+
+ t = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
+
+ if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
+ {
+ cp_lexer_consume_token (parser->lexer);
+
+ t = cp_parser_expression (parser);
+ if (t == error_mark_node
+ || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ }
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_ASYNC, "async", loc);
+
+ c = build_omp_clause (loc, OMP_CLAUSE_ASYNC);
+ OMP_CLAUSE_ASYNC_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ 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
+cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
+ const char *where, cp_token *pragma_tok,
+ bool finish_p = true)
+{
+ tree clauses = NULL;
+ bool first = true;
+
+ while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+ {
+ location_t here;
+ pragma_omp_clause c_kind;
+ const char *c_name;
+ tree prev = clauses;
+
+ if (!first && cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+
+ here = cp_lexer_peek_token (parser->lexer)->location;
+ c_kind = cp_parser_omp_clause_name (parser);
+
+ switch (c_kind)
+ {
+ case PRAGMA_OACC_CLAUSE_ASYNC:
+ clauses = cp_parser_oacc_clause_async (parser, clauses);
+ c_name = "async";
+ break;
+ case PRAGMA_OMP_CLAUSE_COLLAPSE:
+ clauses = cp_parser_omp_clause_collapse (parser, clauses, here);
+ c_name = "collapse";
+ break;
+ case PRAGMA_OACC_CLAUSE_COPY:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyin";
+ break;
+ case PRAGMA_OACC_CLAUSE_COPYOUT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyout";
+ break;
+ case PRAGMA_OACC_CLAUSE_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "create";
+ break;
+ case PRAGMA_OACC_CLAUSE_DELETE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "delete";
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "device";
+ break;
+ case PRAGMA_OACC_CLAUSE_DEVICEPTR:
+ clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
+ c_name = "deviceptr";
+ break;
+ case PRAGMA_OMP_CLAUSE_IF:
+ clauses = cp_parser_omp_clause_if (parser, clauses, here);
+ c_name = "if";
+ break;
+ case PRAGMA_OACC_CLAUSE_NUM_GANGS:
+ clauses = cp_parser_omp_clause_num_gangs (parser, clauses);
+ c_name = "num_gangs";
+ break;
+ case PRAGMA_OACC_CLAUSE_NUM_WORKERS:
+ clauses = cp_parser_omp_clause_num_workers (parser, clauses);
+ c_name = "num_workers";
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present";
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copy";
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyin";
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyout";
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_create";
+ break;
+ case PRAGMA_OMP_CLAUSE_REDUCTION:
+ clauses = cp_parser_omp_clause_reduction (parser, clauses);
+ c_name = "reduction";
+ break;
+ case PRAGMA_OACC_CLAUSE_SELF:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "self";
+ break;
+ case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
+ clauses =
+ cp_parser_oacc_clause_vector_length (parser, clauses);
+ c_name = "vector_length";
+ break;
+ case PRAGMA_OACC_CLAUSE_WAIT:
+ clauses = cp_parser_oacc_clause_wait (parser, clauses);
+ c_name = "wait";
+ break;
+ default:
+ cp_parser_error (parser, "expected clause");
+ goto saw_error;
+ }
+
+ first = false;
+
+ if (((mask >> c_kind) & 1) == 0)
+ {
+ /* 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:
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+
+ if (finish_p)
+ return 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. */
@@ -30842,6 +31376,283 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
return true;
}
+/* OpenACC 2.0:
+ # pragma acc cache (variable-list) new-line
+*/
+
+static tree
+cp_parser_oacc_cache (cp_parser *parser)
+{
+ cp_parser_omp_var_list (parser, OMP_NO_CLAUSE_CACHE, NULL_TREE);
+ cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
+
+ return NULL_TREE;
+}
+
+/* OpenACC 2.0:
+ # pragma acc data oacc-data-clause[optseq] new-line
+ structured-block */
+
+#define OACC_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+
+static tree
+cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+ "#pragma acc data", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_data (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
+*/
+
+#define OACC_ENTER_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+#define OACC_EXIT_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+static tree
+cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
+ bool enter)
+{
+ location_t loc = pragma_tok->location;
+ tree stmt, clauses;
+ const char *p = "";
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ p = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
+
+ if (strcmp (p, "data") != 0)
+ {
+ error_at (loc, "expected %<data%> after %<#pragma acc %s%>",
+ enter ? "enter" : "exit");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return NULL_TREE;
+ }
+
+ cp_lexer_consume_token (parser->lexer);
+
+ if (enter)
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK,
+ "#pragma acc enter data", pragma_tok);
+ else
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_EXIT_DATA_CLAUSE_MASK,
+ "#pragma acc exit data", pragma_tok);
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (loc, "%<#pragma acc %s data%> has no data movement clause",
+ enter ? "enter" : "exit");
+ return NULL_TREE;
+ }
+
+ 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);
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc kernels oacc-kernels-clause[optseq] new-line
+ structured-block */
+
+#define OACC_KERNELS_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
+
+static tree
+cp_parser_oacc_kernels (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
+ "#pragma acc kernels", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_kernels (clauses, block);
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc loop oacc-loop-clause[optseq] new-line
+ structured-block */
+
+#define OACC_LOOP_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION))
+
+static tree
+cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK,
+ "#pragma acc loop", pragma_tok);
+
+ block = begin_omp_structured_block ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ stmt = cp_parser_omp_for_loop (parser, OACC_LOOP, clauses, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ add_stmt (finish_omp_structured_block (block));
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc parallel oacc-parallel-clause[optseq] new-line
+ structured-block */
+
+#define OACC_PARALLEL_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
+
+static tree
+cp_parser_oacc_parallel (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
+ "#pragma acc parallel", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_parallel (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_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
+
+static tree
+cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK,
+ "#pragma acc update", pragma_tok);
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (pragma_tok->location,
+ "%<#pragma acc update%> must contain at least one: "
+ "%<device%>, %<host%> or %<self%> clause");
+ return NULL_TREE;
+ }
+
+ stmt = make_node (OACC_UPDATE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_UPDATE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc wait [(intseq)] oacc-wait-clause[optseq] new-line
+*/
+
+#define OACC_WAIT_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC))
+
+static tree
+cp_parser_oacc_wait (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree clauses, list = NULL_TREE, stmt = NULL_TREE;
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+
+ if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
+ list = cp_parser_oacc_wait_list (parser, loc, list);
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_WAIT_CLAUSE_MASK,
+ "#pragma acc wait", pragma_tok);
+
+ stmt = c_finish_oacc_wait (loc, list, clauses);
+
+ return stmt;
+}
+
/* OpenMP 4.0:
# pragma omp declare simd declare-simd-clauses[optseq] new-line */
@@ -31516,6 +32327,33 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
switch (pragma_tok->pragma_kind)
{
+ case PRAGMA_OACC_CACHE:
+ stmt = cp_parser_oacc_cache (parser);
+ break;
+ case PRAGMA_OACC_DATA:
+ stmt = cp_parser_oacc_data (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_ENTER_DATA:
+ stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, true);
+ break;
+ case PRAGMA_OACC_EXIT_DATA:
+ stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, false);
+ break;
+ case PRAGMA_OACC_KERNELS:
+ stmt = cp_parser_oacc_kernels (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_LOOP:
+ stmt = cp_parser_oacc_loop (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_PARALLEL:
+ stmt = cp_parser_oacc_parallel (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_UPDATE:
+ stmt = cp_parser_oacc_update (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_WAIT:
+ stmt = cp_parser_oacc_wait (parser, pragma_tok);
+ break;
case PRAGMA_OMP_ATOMIC:
cp_parser_omp_atomic (parser, pragma_tok);
return;
@@ -32058,6 +32896,15 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
cp_parser_omp_declare (parser, pragma_tok, context);
return false;
+ case PRAGMA_OACC_CACHE:
+ case PRAGMA_OACC_DATA:
+ case PRAGMA_OACC_ENTER_DATA:
+ case PRAGMA_OACC_EXIT_DATA:
+ case PRAGMA_OACC_KERNELS:
+ case PRAGMA_OACC_PARALLEL:
+ case PRAGMA_OACC_LOOP:
+ case PRAGMA_OACC_UPDATE:
+ case PRAGMA_OACC_WAIT:
case PRAGMA_OMP_ATOMIC:
case PRAGMA_OMP_CRITICAL:
case PRAGMA_OMP_DISTRIBUTE:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 639702a..768ff4d 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5536,6 +5536,44 @@ finish_omp_clauses (tree clauses)
}
break;
+ case OMP_CLAUSE_ASYNC:
+ t = OMP_CLAUSE_ASYNC_EXPR (c);
+ if (t == error_mark_node)
+ remove = true;
+ else if (!type_dependent_expression_p (t)
+ && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ error ("%<async%> expression must be integral");
+ remove = true;
+ }
+ else
+ {
+ t = mark_rvalue_use (t);
+ if (!processing_template_decl)
+ t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+ OMP_CLAUSE_ASYNC_EXPR (c) = t;
+ }
+ break;
+
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ t = OMP_CLAUSE_VECTOR_LENGTH_EXPR (c);
+ t = maybe_convert_cond (t);
+ if (t == error_mark_node)
+ remove = true;
+ else if (!processing_template_decl)
+ t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t;
+ break;
+
+ case OMP_CLAUSE_WAIT:
+ t = OMP_CLAUSE_WAIT_EXPR (c);
+ if (t == error_mark_node)
+ remove = true;
+ else if (!processing_template_decl)
+ t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+ OMP_CLAUSE_WAIT_EXPR (c) = t;
+ break;
+
case OMP_CLAUSE_THREAD_LIMIT:
t = OMP_CLAUSE_THREAD_LIMIT_EXPR (c);
if (t == error_mark_node)
@@ -6053,6 +6091,60 @@ finish_omp_structured_block (tree block)
return do_poplevel (block);
}
+/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_DATA. */
+
+tree
+finish_oacc_data (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DATA_CLAUSES (stmt) = clauses;
+ OACC_DATA_BODY (stmt) = block;
+
+ return add_stmt (stmt);
+}
+
+/* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_KERNELS. */
+
+tree
+finish_oacc_kernels (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_KERNELS);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_KERNELS_CLAUSES (stmt) = clauses;
+ OACC_KERNELS_BODY (stmt) = block;
+
+ return add_stmt (stmt);
+}
+
+/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_PARALLEL. */
+
+tree
+finish_oacc_parallel (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_PARALLEL);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_PARALLEL_CLAUSES (stmt) = clauses;
+ OACC_PARALLEL_BODY (stmt) = block;
+
+ return add_stmt (stmt);
+}
+
/* Similarly, except force the retention of the BLOCK. */
tree
