OpenACC host_data support.

	gcc/
	* gimple-pretty-print.c (dump_gimple_omp_target): Add host_data
	support.
	* gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_HOST_DATA.
	(is_gimple_omp_oacc): Add support for above.
	* gimplify.c (omp_region_type): Add ORT_ACC_HOST_DATA.
	(omp_notice_variable): Diagnose undefined implicit uses of
	use_device variables in offloaded regions.
	(gimplify_scan_omp_clauses): Add host_data, use_device
	support. Diagnose undefined mapping of use_device variables in
	OpenACC clauses.
	(gimplify_omp_workshare): Add host_data support.
	(gimplify_expr): Likewise.
	* omp-builtins.def (BUILT_IN_GOACC_HOST_DATA): New.
	* omp-low.c (lookup_decl_in_outer_ctx)
	(maybe_lookup_decl_in_outer_ctx): Add optional argument to skip
	host_data regions.
	(scan_sharing_clauses): Support use_device.
	(check_omp_nesting_restrictions): Support host_data.
	(expand_omp_target): Support host_data.
	(lower_omp_target): Skip over outer host_data regions when looking
	up decls. Support use_device.
	(make_gimple_omp_edges): Support host_data.
	* tree-nested.c (convert_nonlocal_omp_clauses): Add use_device
	clause.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Add use_device support.
	(c_parser_oacc_clause_use_device): New function.
	(c_parser_oacc_all_clauses): Add use_device support.
	(OACC_HOST_DATA_CLAUSE_MASK): New macro.
	(c_parser_oacc_host_data): New function.
	(c_parser_omp_construct): Add host_data support.
	* c-tree.h (c_finish_oacc_host_data): Add prototype.
	* c-typeck.c (c_finish_oacc_host_data): New function.
	(c_finish_omp_clauses): Add use_device support.

	gcc/cp/
	* cp-tree.h (finish_oacc_host_data): Add prototype.
	* parser.c (cp_parser_omp_clause_name): Add use_device support.
	(cp_parser_oacc_all_clauses): Add use_device support.
	(OACC_HOST_DATA_CLAUSE_MASK): New macro.
	(cp_parser_oacc_host_data): New function.
	(cp_parser_omp_construct): Add host_data support.
	(cp_parser_pragma): Add host_data support.
	* semantics.c (finish_omp_clauses): Add use_device support.
	(finish_oacc_host_data): New function.

	gcc/c-family/
	* c-pragma.c (oacc_pragmas): Add PRAGMA_OACC_HOST_DATA.
	* c-pragma.h (pragma_kind): Add PRAGMA_OACC_HOST_DATA.
	(pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_USE_DEVICE.

	libgomp/
	* oacc-parallel.c (GOACC_host_data): New function.
	* libgomp.map (GOACC_host_data): Add to GOACC_2.0.1.
	* testsuite/libgomp.oacc-c-c++-common/host_data-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-5.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.


Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: James Norris <James_Norris@mentor.com>

From-SVN: r231118

4 files changed, 69 insertions, 0 deletions

diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index b424336..9cc0a53 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,17 @@
+2015-12-01  Julian Brown  <julian@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <James_Norris@mentor.com>
+
+	* cp-tree.h (finish_oacc_host_data): Add prototype.
+	* parser.c (cp_parser_omp_clause_name): Add use_device support.
+	(cp_parser_oacc_all_clauses): Add use_device support.
+	(OACC_HOST_DATA_CLAUSE_MASK): New macro.
+	(cp_parser_oacc_host_data): New function.
+	(cp_parser_omp_construct): Add host_data support.
+	(cp_parser_pragma): Add host_data support.
+	* semantics.c (finish_omp_clauses): Add use_device support.
+	(finish_oacc_host_data): New function.
+
 2015-11-27  Martin Liska  <mliska@suse.cz>
 
 	PR c++/68312
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index caa601d..38ae70f 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -6360,6 +6360,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_omp_construct		(enum tree_code, tree, tree);
 extern tree begin_omp_parallel			(void);
 extern tree finish_omp_parallel			(tree, tree);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 90a0673..f78df02 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29232,6 +29232,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", p))
@@ -31598,6 +31600,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_SEQ:
 	  clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
 						 clauses, here);
@@ -34509,6 +34516,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 declare oacc-data-clause[optseq] new-line
 */
@@ -36068,6 +36099,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:
     case PRAGMA_OACC_PARALLEL:
       strcpy (p_name, "#pragma acc");
@@ -36645,6 +36679,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:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index e7e5d8e..3bb6184 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6835,6 +6835,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	    }
 	  break;
 
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  field_ok = allow_fields;
@@ -7390,6 +7391,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 OMP construct CODE, with BODY and CLAUSES as its compound
    statement.  */