This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[OpenACC 2/7] host_data construct (C FE)
- From: James Norris <jnorris at codesourcery dot com>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Cc: "Joseph S. Myers" <joseph at codesourcery dot com>, Nathan Sidwell <Nathan_Sidwell at mentor dot com>, Jakub Jelinek <jakub at redhat dot com>
- Date: Thu, 22 Oct 2015 14:14:54 -0500
- Subject: [OpenACC 2/7] host_data construct (C FE)
- Authentication-results: sourceware.org; auth=none
- References: <56293476 dot 5020801 at codesourcery dot com>
gcc/c/c-parser.c b/gcc/c/c-parser.c
- Add handling of use_device clause in c_parser_omp_clause_name().
- Add new function c_parser_oacc_clause_use_device() to handle
use_device clause with host_data.
- Add handling of use_device clause in c_parser_oacc_all_clauses().
- Add new macro OACC_HOST_DATA_CLAUSE_MASK.
- Add new function c_parser_oacc_host_data() to handle host_data.
- Add handling of host_data pragma to c_parser_omp_construct().
gcc/c/c-tree.h b/gcc/c/c-tree.h
- Add definition for c_finish_oacc_host_data().
gcc/c/c-typeck.c b/gcc/c/c-typeck.c
- Add new function c_finish_oacc_host_data().
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 704ebc6..ead98b9 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10116,6 +10116,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("use_device", p))
+ result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
break;
case 'v':
if (!strcmp ("vector", p))
@@ -11219,6 +11221,15 @@ c_parser_oacc_clause_async (c_parser *parser, tree list)
return list;
}
+/* OpenACC 2.0:
+ use_device ( variable-list ) */
+
+static tree
+c_parser_oacc_clause_use_device (c_parser *parser, tree list)
+{
+ return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list);
+}
+
/* OpenACC:
wait ( int-expr-list ) */
@@ -12474,6 +12485,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "self";
break;
+ case PRAGMA_OACC_CLAUSE_USE_DEVICE:
+ clauses = c_parser_oacc_clause_use_device (parser, clauses);
+ c_name = "use_device";
+ break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
clauses = c_parser_omp_clause_vector_length (parser, clauses);
c_name = "vector_length";
@@ -13003,6 +13018,29 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
/* OpenACC 2.0:
+ # pragma acc host_data oacc-data-clause[optseq] new-line
+ structured-block
+*/
+
+#define OACC_HOST_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
+
+static tree
+c_parser_oacc_host_data (location_t loc, c_parser *parser)
+{
+ tree stmt, clauses, block;
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK,
+ "#pragma acc host_data");
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+ stmt = c_finish_oacc_host_data (loc, clauses, block);
+ return stmt;
+}
+
+
+/* OpenACC 2.0:
# pragma acc loop oacc-loop-clause[optseq] new-line
structured-block
@@ -16075,6 +16113,9 @@ c_parser_omp_construct (c_parser *parser)
case PRAGMA_OACC_DATA:
stmt = c_parser_oacc_data (loc, parser);
break;
+ case PRAGMA_OACC_HOST_DATA:
+ stmt = c_parser_oacc_host_data (loc, parser);
+ break;
case PRAGMA_OACC_KERNELS:
strcpy (p_name, "#pragma acc");
stmt = c_parser_oacc_kernels (loc, parser, p_name);
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index bee03d3..a9c5975 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -643,6 +643,7 @@ extern tree c_expr_to_decl (tree, bool *, bool *);
extern tree c_finish_oacc_parallel (location_t, tree, tree);
extern tree c_finish_oacc_kernels (location_t, tree, tree);
extern tree c_finish_oacc_data (location_t, tree, tree);
+extern tree c_finish_oacc_host_data (location_t, tree, tree);
extern tree c_begin_omp_parallel (void);
extern tree c_finish_omp_parallel (location_t, tree, tree);
extern tree c_begin_omp_task (void);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index bc43602..a5e2a4a 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -11510,6 +11510,25 @@ c_finish_oacc_data (location_t loc, tree clauses, tree block)
return add_stmt (stmt);
}
+/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_HOST_DATA. */
+
+tree
+c_finish_oacc_host_data (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_HOST_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_HOST_DATA_CLAUSES (stmt) = clauses;
+ OACC_HOST_DATA_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
/* Like c_begin_compound_stmt, except force the retention of the BLOCK. */
tree
@@ -12942,6 +12961,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_USE_DEVICE:
pc = &OMP_CLAUSE_CHAIN (c);
continue;