Initial OpenACC support to C++ front-end: parallel, kernels and data
construct with data clauses.
gcc/cp/
* cp-tree.h (finish_oacc_data): New function prototype.
(finish_oacc_kernels, finish_oacc_parallel): Likewise.
* parser.c (cp_parser_omp_clause_name): Support data clauses.
(cp_parser_oacc_data_clause): New function.
(cp_parser_oacc_data_clause_deviceptr, cp_parser_oacc_all_clauses,
cp_parser_oacc_data, cp_parser_oacc_kernels,
cp_parser_oacc_parallel): Likewise.
(OACC_DATA_CLAUSE_MASK): New define.
(OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Likewise.
(cp_parser_omp_construct, cp_parser_pragma): Support OpenACC
directives.
* semantics.c (finish_oacc_data): New function.
(finish_oacc_kernels, finish_oacc_parallel): Likewise.
>From 422e398c375ce6a62691bfc4cc9dd71c38f87b40 Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usma...@samsung.com>
Date: Fri, 7 Mar 2014 15:09:47 +0400
Subject: [PATCH 1/2] Initial OpenACC support to C++ front-end
---
gcc/cp/cp-tree.h | 3 +
gcc/cp/parser.c | 347 ++++++++++++++++++++++++++++++++++++++++++++++++++++-
gcc/cp/semantics.c | 54 +++++++++
3 files changed, 403 insertions(+), 1 deletion(-)
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 7681b27..3d734ac 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5822,6 +5822,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 ff481ab..ba9cda65 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -26868,16 +26868,26 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'c':
if (!strcmp ("collapse", p))
result = PRAGMA_OMP_CLAUSE_COLLAPSE;
+ else if (!strcmp ("copy", p))
+ result = PRAGMA_OMP_CLAUSE_COPY;
else if (!strcmp ("copyin", p))
result = PRAGMA_OMP_CLAUSE_COPYIN;
+ else if (!strcmp ("copyout", p))
+ result = PRAGMA_OMP_CLAUSE_COPYOUT;
else if (!strcmp ("copyprivate", p))
result = PRAGMA_OMP_CLAUSE_COPYPRIVATE;
+ else if (!strcmp ("create", p))
+ result = PRAGMA_OMP_CLAUSE_CREATE;
break;
case 'd':
- if (!strcmp ("depend", p))
+ if (!strcmp ("delete", p))
+ result = PRAGMA_OMP_CLAUSE_DELETE;
+ else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
+ else if (!strcmp ("deviceptr", p))
+ result = PRAGMA_OMP_CLAUSE_DEVICEPTR;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -26926,6 +26936,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;
@@ -27133,6 +27159,111 @@ 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_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;
+ }
+ 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; 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;
+}
+
/* OpenMP 3.0:
collapse ( constant-expression ) */
@@ -28045,6 +28176,101 @@ cp_parser_omp_clause_proc_bind (cp_parser *parser, tree list,
return list;
}
+/* Parse all OpenACC clauses. The set clauses allowed by the directive
+ is a bitmask in MASK. Return the list of clauses found. */
+
+static tree
+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_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_DEVICEPTR:
+ clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
+ c_name = "deviceptr";
+ 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;
+ 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. */
@@ -28341,6 +28567,113 @@ cp_parser_omp_structured_block (cp_parser *parser)
return finish_omp_structured_block (stmt);
}
+/* OpenACC 2.0:
+ # pragma acc data oacc-data-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_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 kernels oacc-kernels-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_KERNELS_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_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_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_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 parallel oacc-parallel-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_PARALLEL_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_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_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_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;
+}
/* OpenMP 2.5:
# pragma omp atomic new-line
expression-stmt
@@ -30903,6 +31236,15 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
switch (pragma_tok->pragma_kind)
{
+ case PRAGMA_OACC_DATA:
+ stmt = cp_parser_oacc_data (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_KERNELS:
+ stmt = cp_parser_oacc_kernels (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_PARALLEL:
+ stmt = cp_parser_oacc_parallel (parser, pragma_tok);
+ break;
case PRAGMA_OMP_ATOMIC:
cp_parser_omp_atomic (parser, pragma_tok);
return;
@@ -31415,6 +31757,9 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
cp_parser_omp_declare (parser, pragma_tok, context);
return false;
+ case PRAGMA_OACC_DATA:
+ case PRAGMA_OACC_KERNELS:
+ case PRAGMA_OACC_PARALLEL:
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 9fb4fc0..5b9b879 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5986,6 +5986,60 @@ finish_omp_structured_block (tree block)
return do_poplevel (block);
}
+/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_PARALLEL. */
+
+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_PARALLEL, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_PARALLEL. */
+
+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
--
1.8.3.2
