summaryrefslogtreecommitdiff
path: root/gcc/cp
diff options
context:
space:
mode:
Diffstat (limited to 'gcc/cp')
-rw-r--r--gcc/cp/ChangeLog45
-rw-r--r--gcc/cp/cp-tree.h3
-rw-r--r--gcc/cp/parser.c880
-rw-r--r--gcc/cp/semantics.c102
4 files changed, 1019 insertions, 11 deletions
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 02b4fac907b..543f4d9a4d9 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,48 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+ Jakub Jelinek <jakub@redhat.com>
+
+ * parser.c: Include "gomp-constants.h".
+ (cp_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
+ omp_clause_map_kind. Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
+ Use OMP_CLAUSE_SET_MAP_KIND.
+ (cp_parser_omp_construct, cp_parser_pragma): Handle
+ PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA,
+ PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_LOOP, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
+ (cp_parser_omp_clause_name): Handle "async", "copy", "copyout",
+ "create", "delete", "deviceptr", "host", "num_gangs",
+ "num_workers", "present", "present_or_copy", "pcopy",
+ "present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
+ "present_or_create", "pcreate", "vector_length", "wait".
+ (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
+ (OACC_EXIT_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
+ (OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
+ (OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
+ (cp_parser_omp_var_list_no_open): Handle OMP_CLAUSE__CACHE_.
+ (cp_parser_oacc_data_clause, cp_parser_oacc_data_clause_deviceptr)
+ (cp_parser_oacc_clause_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-tree.h (finish_oacc_data, finish_oacc_kernels)
+ (finish_oacc_parallel): New prototypes.
+ * semantics.c: Include "gomp-constants.h".
+ (handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR. Use
+ GOMP_MAP_* instead of OMP_CLAUSE_MAP_*. Use
+ OMP_CLAUSE_SET_MAP_KIND.
+ (finish_omp_clauses): Handle OMP_CLAUSE_ASYNC,
+ OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_WAIT, OMP_CLAUSE__CACHE_.
+ Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
+ (finish_oacc_data, finish_oacc_kernels, finish_oacc_parallel): New
+ functions.
+
2015-01-14 Paolo Carlini <paolo.carlini@oracle.com>
PR c++/58671
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 77f2b5b918c..10c63fd1b4d 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5990,6 +5990,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 3290dfa7e59..bfa3d81bf74 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -60,6 +60,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. */
@@ -27542,6 +27543,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))
@@ -27556,20 +27559,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;
@@ -27581,6 +27594,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_HOST;
+ break;
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
@@ -27606,10 +27623,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))
@@ -27618,6 +27639,20 @@ 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 ("proc_bind", p))
result = PRAGMA_OMP_CLAUSE_PROC_BIND;
break;
@@ -27632,6 +27667,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))
@@ -27652,9 +27689,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;
}
}
@@ -27730,6 +27773,14 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
{
switch (kind)
{
+ case OMP_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:
@@ -27760,6 +27811,26 @@ 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_CLAUSE__CACHE_)
+ {
+ 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;
@@ -27822,6 +27893,222 @@ 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 gomp_map_kind kind;
+ switch (c_kind)
+ {
+ case PRAGMA_OACC_CLAUSE_COPY:
+ kind = GOMP_MAP_FORCE_TOFROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_COPYIN:
+ kind = GOMP_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OACC_CLAUSE_COPYOUT:
+ kind = GOMP_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_CREATE:
+ kind = GOMP_MAP_FORCE_ALLOC;
+ break;
+ case PRAGMA_OACC_CLAUSE_DELETE:
+ kind = GOMP_MAP_FORCE_DEALLOC;
+ break;
+ case PRAGMA_OACC_CLAUSE_DEVICE:
+ kind = GOMP_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OACC_CLAUSE_HOST:
+ case PRAGMA_OACC_CLAUSE_SELF:
+ kind = GOMP_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT:
+ kind = GOMP_MAP_FORCE_PRESENT;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
+ kind = GOMP_MAP_TOFROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
+ kind = GOMP_MAP_TO;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ kind = GOMP_MAP_FROM;
+ break;
+ case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
+ kind = GOMP_MAP_ALLOC;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ 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_SET_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_OACC_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_SET_MAP_KIND (u, GOMP_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;
+
+ 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 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 ) */
@@ -28010,6 +28297,42 @@ 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;
+
+ 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 ) */
@@ -28040,6 +28363,43 @@ 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;
+
+ 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_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 */
@@ -28562,7 +28922,7 @@ static tree
cp_parser_omp_clause_map (cp_parser *parser, tree list)
{
tree nlist, c;
- enum omp_clause_map_kind kind = OMP_CLAUSE_MAP_TOFROM;
+ enum gomp_map_kind kind = GOMP_MAP_TOFROM;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
@@ -28574,13 +28934,13 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
const char *p = IDENTIFIER_POINTER (id);
if (strcmp ("alloc", p) == 0)
- kind = OMP_CLAUSE_MAP_ALLOC;
+ kind = GOMP_MAP_ALLOC;
else if (strcmp ("to", p) == 0)
- kind = OMP_CLAUSE_MAP_TO;
+ kind = GOMP_MAP_TO;
else if (strcmp ("from", p) == 0)
- kind = OMP_CLAUSE_MAP_FROM;
+ kind = GOMP_MAP_FROM;
else if (strcmp ("tofrom", p) == 0)
- kind = OMP_CLAUSE_MAP_TOFROM;
+ kind = GOMP_MAP_TOFROM;
else
{
cp_parser_error (parser, "invalid map kind");
@@ -28597,7 +28957,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
NULL);
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_MAP_KIND (c) = kind;
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
return nlist;
}
@@ -28734,6 +29094,178 @@ 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_OACC_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_OACC_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_OACC_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_OACC_CLAUSE_HOST:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "host";
+ break;
+ case PRAGMA_OACC_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_OACC_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 %<#pragma acc%> 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. */
@@ -30953,6 +31485,304 @@ 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)
+{
+ tree stmt, clauses;
+
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
+ clauses = finish_omp_clauses (clauses);
+
+ cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
+
+ stmt = make_node (OACC_CACHE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_CACHE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+
+ return stmt;
+}
+
+/* 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_OACC_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_OACC_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
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_ENTER_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_OACC_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)
+{
+ 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_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_OACC_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_OACC_CLAUSE_COLLAPSE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_OACC_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_OACC_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_OACC_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_OACC_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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%> 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_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 */
@@ -31627,6 +32457,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;
@@ -32169,6 +33026,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 15b8d0111a0..915048daf0a 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -66,6 +66,7 @@ along with GCC; see the file COPYING3. If not see
#include "omp-low.h"
#include "builtins.h"
#include "convert.h"
+#include "gomp-constants.h"
/* There routines provide a modular interface to perform many parsing
operations. They may therefore be used during actual parsing, or
@@ -4670,7 +4671,7 @@ handle_omp_array_sections (tree c)
return false;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
if (!cxx_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
@@ -4694,7 +4695,7 @@ handle_omp_array_sections (tree c)
{
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_MAP_KIND (c3) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_POINTER);
OMP_CLAUSE_DECL (c3) = ptr;
OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
OMP_CLAUSE_SIZE (c3) = size_zero_node;
@@ -5571,6 +5572,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)
@@ -5721,6 +5760,7 @@ finish_omp_clauses (tree clauses)
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
@@ -5749,7 +5789,7 @@ finish_omp_clauses (tree clauses)
if (processing_template_decl)
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
break;
if (DECL_P (t))
error ("%qD is not a variable in %qs clause", t,
@@ -5770,7 +5810,7 @@ finish_omp_clauses (tree clauses)
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
&& !type_dependent_expression_p (t)
&& !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
== REFERENCE_TYPE)
@@ -6088,6 +6128,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
generated by cgit v1.2.3 (git 2.25.1) at 2024-12-02 04:32:48 +0000