@@ -5187,6 +5187,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
@@ -5205,6 +5210,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
};
@@ -5297,6 +5304,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]);
@@ -5318,6 +5333,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;
@@ -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);
@@ -1186,6 +1186,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");
@@ -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
@@ -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 },
+};
static const struct omp_pragma_def omp_pragmas[] = {
{ "atomic", PRAGMA_OMP_ATOMIC },
{ "barrier", PRAGMA_OMP_BARRIER },
@@ -1212,11 +1222,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 +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);
@@ -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,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,
/* Clauses for Cilk Plus SIMD-enabled function. */
PRAGMA_CILK_CLAUSE_NOMASK,
@@ -1192,6 +1192,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)
@@ -5923,6 +5923,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);
@@ -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;
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))
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;
@@ -27476,6 +27488,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_OMP_CLAUSE_HOST;
+ break;
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
@@ -27501,10 +27517,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_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;
break;
case 'o':
if (!strcmp ("ordered", p))
@@ -27513,6 +27533,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_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))
result = PRAGMA_OMP_CLAUSE_PROC_BIND;
break;
@@ -27527,6 +27563,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_OMP_CLAUSE_SELF;
else if (!strcmp ("shared", p))
result = PRAGMA_OMP_CLAUSE_SHARED;
else if (!strcmp ("simdlen", p))
@@ -27547,9 +27585,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_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;
}
}
@@ -27625,6 +27669,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:
@@ -27655,6 +27707,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;
@@ -27717,6 +27792,233 @@ 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_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 = 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);
+
+ /* 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:
+ 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))
+ || !tree_fits_shwi_p (t)
+ || (n = tree_to_shwi (t)) <= 0
+ || (int) n != n)
+ {
+ 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 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");
+ targ = error_mark_node;
+ }
+ 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 ) */
@@ -27905,6 +28207,46 @@ 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))
+ || !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_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 ) */
@@ -27935,6 +28277,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 */
@@ -28629,6 +29012,180 @@ 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;
+
+ /* TODO XXX: FIX -1 (acc_async_noval). */
+ t = build_int_cst (integer_type_node, -1);
+
+ 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_OMP_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_OMP_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_OMP_CLAUSE_COPYOUT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "create";
+ break;
+ case PRAGMA_OMP_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_OMP_CLAUSE_DEVICEPTR:
+ clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
+ c_name = "deviceptr";
+ break;
+ case PRAGMA_OMP_CLAUSE_HOST:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "host";
+ break;
+ case PRAGMA_OMP_CLAUSE_IF:
+ clauses = cp_parser_omp_clause_if (parser, clauses, here);
+ c_name = "if";
+ break;
+ case PRAGMA_OMP_CLAUSE_NUM_GANGS:
+ clauses = cp_parser_omp_clause_num_gangs (parser, clauses);
+ c_name = "num_gangs";
+ break;
+ case PRAGMA_OMP_CLAUSE_NUM_WORKERS:
+ clauses = cp_parser_omp_clause_num_workers (parser, clauses);
+ c_name = "num_workers";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyout";
+ break;
+ case PRAGMA_OMP_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_OMP_CLAUSE_SELF:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "self";
+ break;
+ case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH:
+ clauses =
+ cp_parser_oacc_clause_vector_length (parser, clauses);
+ c_name = "vector_length";
+ break;
+ case PRAGMA_OMP_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. */
@@ -30848,6 +31405,295 @@ 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_token *pragma_tok __attribute__((unused)))
+{
+ 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_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
+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
+
+ 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 tree
+cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
+ bool enter)
+{
+ tree stmt, clauses;
+
+ 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%>");
+ 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");
+ 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 (pragma_tok->location,
+ "%<#pragma acc enter data%> has no data movement clause");
+ 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, pragma_tok->location);
+ 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_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
+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_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
+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_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 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%> or %<host/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
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_WAIT_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_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 */
@@ -31522,6 +32368,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, pragma_tok);
+ 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;
@@ -32064,6 +32937,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:
@@ -5516,6 +5516,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)
@@ -6033,6 +6071,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