@@ -6318,6 +6318,7 @@ 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_host_data (tree, tree);
extern tree finish_oacc_kernels (tree, tree);
extern tree finish_oacc_parallel (tree, tree);
extern tree begin_omp_parallel (void);
@@ -29235,6 +29235,8 @@ cp_parser_omp_clause_name (cp_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_length", p))
@@ -31381,6 +31383,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "self";
break;
+ case PRAGMA_OACC_CLAUSE_USE_DEVICE:
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE,
+ clauses);
+ c_name = "use_device";
+ break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
clauses = cp_parser_oacc_clause_vector_length (parser, clauses);
c_name = "vector_length";
@@ -34221,6 +34228,30 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
return stmt;
}
+#define OACC_HOST_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
+
+/* OpenACC 2.0:
+ # pragma acc host_data <clauses> new-line
+ structured-block */
+
+static tree
+cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK,
+ "#pragma acc host_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_host_data (clauses, block);
+ return stmt;
+}
+
/* OpenACC 2.0:
# pragma acc enter data oacc-enter-data-clause[optseq] new-line
@@ -35288,6 +35319,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
case PRAGMA_OACC_EXIT_DATA:
stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, false);
break;
+ case PRAGMA_OACC_HOST_DATA:
+ stmt = cp_parser_oacc_host_data (parser, pragma_tok);
+ break;
case PRAGMA_OACC_KERNELS:
stmt = cp_parser_oacc_kernels (parser, pragma_tok);
break;
@@ -35856,6 +35890,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
case PRAGMA_OACC_DATA:
case PRAGMA_OACC_ENTER_DATA:
case PRAGMA_OACC_EXIT_DATA:
+ case PRAGMA_OACC_HOST_DATA:
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_LOOP:
@@ -6689,6 +6689,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ case OMP_CLAUSE_USE_DEVICE:
break;
case OMP_CLAUSE_INBRANCH:
@@ -7119,6 +7120,24 @@ finish_oacc_data (tree clauses, tree block)
return add_stmt (stmt);
}
+/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound
+ statement. */
+
+tree
+finish_oacc_host_data (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_HOST_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_HOST_DATA_CLAUSES (stmt) = clauses;
+ OACC_HOST_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. */