diff mbox

[OpenACC,0/7] host_data construct

Message ID 20151130193034.71143aa4@octopus
State New
Headers show

Commit Message

Julian Brown Nov. 30, 2015, 7:30 p.m. UTC
On Thu, 19 Nov 2015 16:57:23 +0100
Jakub Jelinek <jakub@redhat.com> wrote:

> If it is unclear, I think disallowing acc {parallel,kernels} inside of
> acc host_data might be too big hammer, but perhaps just erroring out
> or warning during gimplification that if you (explicitly or
> implicitly) try to map a var that is in use_device clause in some
> outer context, it is either wrong, unsupported or will not do what
> users think?

I think we can only assume that trying to map a variable declared in
a surrounding use_device clause is undefined behaviour. I haven't had
any response to my questions about host_data & deviceptr on the OpenACC
list.

> > #pragma acc host_data use_device(x)
> > {
> >   target_primitive(x);
> >   #pragma acc parallel deviceptr(x)
> >   {
> >     ...
> >   }
> > }
> 
> Is deviceptr as above meant to work?  That is the OpenACC counterpart
> of is_device_ptr, right?  If yes, then I'd suggest just warning if you
> try to implicitly or explicitly map something use_device in outer
> contexts, and just make sure you don't ICE on the cases where you
> warn. If the standard does not say what it means, then it is
> unspecified behavior...

A problem with deviceptr, unlike is_device_ptr, is that it turns out to
be defined only to work with pointers, not arrays (OpenACC 2.0a
2.6.5.2), and there are no rules describing the latter decaying to the
former. So at least if 'x' is an array, it appears the answer is "no".

So, the attached patch disallows (via raising an error):

* Variables being declared in explicit mapping clauses that are
  declared in enclosing host_data regions.

* Variables being implicitly used (mapped) in offloaded regions that
  are declared in enclosing host_data regions.

It's otherwise equivalent to the previously-posted version, but without
the hacks to {maybe_,}lookup_decl_in_outer_ctx. I added checks for the
above conditions during gimplification, which seemed to be about the
same phase that other similar kinds of errors are diagnosed.

Tests look OK (libgomp/gcc/g++/libstdc++), and the new ones pass.

OK for mainline?

Thanks,

Julian

ChangeLog

    Julian Brown  <julian@codesourcery.com>
    Cesar Philippidis  <cesar@codesourcery.com>
    James Norris  <James_Norris@mentor.com>

    gcc/
    * c-family/c-pragma.c (oacc_pragmas): Add PRAGMA_OACC_HOST_DATA.
    * c-family/c-pragma.h (pragma_kind): Add PRAGMA_OACC_HOST_DATA.
    (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_USE_DEVICE.
    * 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/c-tree.h (c_finish_oacc_host_data): Add prototype.
    * c/c-typeck.c (c_finish_oacc_host_data): New function.
    (c_finish_omp_clauses): Add use_device support.
    * cp/cp-tree.h (finish_oacc_host_data): Add prototype.
    * cp/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.
    * cp/semantics.c (finish_omp_clauses): Add use_device support.
    (finish_oacc_host_data): New function.
    * 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.

    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.

Comments

Jakub Jelinek Dec. 1, 2015, 8:30 a.m. UTC | #1
On Mon, Nov 30, 2015 at 07:30:34PM +0000, Julian Brown wrote:
>     Julian Brown  <julian@codesourcery.com>
>     Cesar Philippidis  <cesar@codesourcery.com>
>     James Norris  <James_Norris@mentor.com>
> 
>     gcc/
>     * c-family/c-pragma.c (oacc_pragmas): Add PRAGMA_OACC_HOST_DATA.
>     * c-family/c-pragma.h (pragma_kind): Add PRAGMA_OACC_HOST_DATA.

c-family/, c/ and cp/ subdirectories have their own ChangeLog, so you need
to split the entry into multiple ChangeLog files and remove the directory
prefixes.

> @@ -6120,6 +6121,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
>  					 (splay_tree_key) decl);
>  		  if (n2)
>  		    {
> +		      if (octx->region_type == ORT_ACC_HOST_DATA)
> +		        error ("variable %qE declared in enclosing "
> +			       "host_data region", DECL_NAME (decl));

%<host_data%> instead?
>  		      nflags |= GOVD_MAP;
>  		      goto found_outer;
>  		    }
> @@ -6418,6 +6422,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>        case OMP_TARGET_DATA:
>        case OMP_TARGET_ENTER_DATA:
>        case OMP_TARGET_EXIT_DATA:
> +      case OACC_HOST_DATA:
>  	ctx->target_firstprivatize_array_bases = true;
>        default:
>  	break;
> @@ -6683,6 +6688,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  	    case OMP_TARGET_DATA:
>  	    case OMP_TARGET_ENTER_DATA:
>  	    case OMP_TARGET_EXIT_DATA:
> +	    case OACC_HOST_DATA:
>  	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
>  		  || (OMP_CLAUSE_MAP_KIND (c)
>  		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
> @@ -6695,6 +6701,22 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  	    }
>  	  if (remove)
>  	    break;
> +	  if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
> +	    {
> +	      struct gimplify_omp_ctx *octx;
> +	      for (octx = outer_ctx; octx; octx = octx->outer_context)
> +	        {
> +		  if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
> +		    break;

Wouldn't it be better to do
		if (octx->region_type != ORT_ACC_HOST_DATA)
		  continue;
here, thus only lookup if you really want to use it?

> +		  splay_tree_node n2
> +		    = splay_tree_lookup (octx->variables,
> +					 (splay_tree_key) decl);
> +		  if (n2 && octx->region_type == ORT_ACC_HOST_DATA)

and remove the && ... part from the condition?

> +		    error_at (OMP_CLAUSE_LOCATION (c), "variable %qE "
> +			      "declared in enclosing host_data region",
> +			      DECL_NAME (decl));
> +		}
> +	    }
>  	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
>  	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
>  				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));

Ok with those changes.

	Jakub
Tom de Vries Dec. 2, 2015, 3:27 p.m. UTC | #2
On 30/11/15 20:30, Julian Brown wrote:
>      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.
>

Hi,

At r231169, I'm seeing these failures for a no-accelerator setup:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/host_data-2.c 
-DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 execution test
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/host_data-4.c 
-DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 execution test
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/host_data-5.c 
-DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 execution test
...

Thanks,
- Tom
diff mbox

Patch

commit ac77cadbe27ca43d036f943c48fb4bf59bc84b36
Author: Julian Brown <julian@codesourcery.com>
Date:   Mon Nov 2 06:31:47 2015 -0800

    OpenACC host_data support using mapping regions.

diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index 12c3e75..56cf697 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1251,6 +1251,7 @@  static const struct omp_pragma_def oacc_pragmas[] = {
   { "declare", PRAGMA_OACC_DECLARE },
   { "enter", PRAGMA_OACC_ENTER_DATA },
   { "exit", PRAGMA_OACC_EXIT_DATA },
+  { "host_data", PRAGMA_OACC_HOST_DATA },
   { "kernels", PRAGMA_OACC_KERNELS },
   { "loop", PRAGMA_OACC_LOOP },
   { "parallel", PRAGMA_OACC_PARALLEL },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 999ac67..dd246b9 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -33,6 +33,7 @@  enum pragma_kind {
   PRAGMA_OACC_DECLARE,
   PRAGMA_OACC_ENTER_DATA,
   PRAGMA_OACC_EXIT_DATA,
+  PRAGMA_OACC_HOST_DATA,
   PRAGMA_OACC_KERNELS,
   PRAGMA_OACC_LOOP,
   PRAGMA_OACC_PARALLEL,
@@ -167,6 +168,7 @@  enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_SELF,
   PRAGMA_OACC_CLAUSE_SEQ,
   PRAGMA_OACC_CLAUSE_TILE,
+  PRAGMA_OACC_CLAUSE_USE_DEVICE,
   PRAGMA_OACC_CLAUSE_VECTOR,
   PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
   PRAGMA_OACC_CLAUSE_WAIT,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 0259f66..d4c512f 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10279,6 +10279,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))
@@ -11631,6 +11633,15 @@  c_parser_oacc_clause_tile (c_parser *parser, tree list)
   return c;
 }
 
+/* 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 ) */
 
@@ -12940,6 +12951,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_SEQ:
 	  clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
 						clauses);
@@ -13590,6 +13605,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
@@ -16897,6 +16935,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:
     case PRAGMA_OACC_PARALLEL:
       strcpy (p_name, "#pragma acc");
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 6bc216a..848131e 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -653,6 +653,7 @@  extern tree c_finish_goto_ptr (location_t, tree);
 extern tree c_expr_to_decl (tree, bool *, bool *);
 extern tree c_finish_omp_construct (location_t, enum tree_code, 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 741c75c..cc2e38e 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -11631,6 +11631,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
@@ -13074,6 +13093,7 @@  c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 	  bitmap_set_bit (&map_head, DECL_UID (t));
 	  goto check_dup_generic;
 
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  t = OMP_CLAUSE_DECL (c);
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.  */
 
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index 7764201..f1abf5c 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1356,6 +1356,9 @@  dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
     case GF_OMP_TARGET_KIND_OACC_DECLARE:
       kind = " oacc_declare";
       break;
+    case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+      kind = " oacc_host_data";
+      break;
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 0b04804..dc61043 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -171,6 +171,7 @@  enum gf_mask {
     GF_OMP_TARGET_KIND_OACC_UPDATE = 8,
     GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9,
     GF_OMP_TARGET_KIND_OACC_DECLARE = 10,
+    GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11,
 
     /* True on an GIMPLE_OMP_RETURN statement if the return does not require
        a thread synchronization via some sort of barrier.  The exact barrier
@@ -6004,6 +6005,7 @@  is_gimple_omp_oacc (const gimple *stmt)
 	case GF_OMP_TARGET_KIND_OACC_UPDATE:
 	case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
 	case GF_OMP_TARGET_KIND_OACC_DECLARE:
+	case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
 	  return true;
 	default:
 	  return false;
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 7fff12f..85f6b1a 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -122,6 +122,7 @@  enum omp_region_type
   ORT_ACC_DATA	= ORT_ACC | ORT_TARGET_DATA, /* Data construct.  */
   ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET,  /* Parallel construct */
   ORT_ACC_KERNELS  = ORT_ACC | ORT_TARGET | 0x80,  /* Kernels construct.  */
+  ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 0x80,  /* Host data.  */
 
   /* Dummy OpenMP region, used to disable expansion of
      DECL_VALUE_EXPRs in taskloop pre body.  */
@@ -6120,6 +6121,9 @@  omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 					 (splay_tree_key) decl);
 		  if (n2)
 		    {
+		      if (octx->region_type == ORT_ACC_HOST_DATA)
+		        error ("variable %qE declared in enclosing "
+			       "host_data region", DECL_NAME (decl));
 		      nflags |= GOVD_MAP;
 		      goto found_outer;
 		    }
@@ -6418,6 +6422,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       case OMP_TARGET_DATA:
       case OMP_TARGET_ENTER_DATA:
       case OMP_TARGET_EXIT_DATA:
+      case OACC_HOST_DATA:
 	ctx->target_firstprivatize_array_bases = true;
       default:
 	break;
@@ -6683,6 +6688,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
+	    case OACC_HOST_DATA:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
 		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
@@ -6695,6 +6701,22 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  if (remove)
 	    break;
+	  if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
+	    {
+	      struct gimplify_omp_ctx *octx;
+	      for (octx = outer_ctx; octx; octx = octx->outer_context)
+	        {
+		  if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
+		    break;
+		  splay_tree_node n2
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) decl);
+		  if (n2 && octx->region_type == ORT_ACC_HOST_DATA)
+		    error_at (OMP_CLAUSE_LOCATION (c), "variable %qE "
+			      "declared in enclosing host_data region",
+			      DECL_NAME (decl));
+		}
+	    }
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -7092,6 +7114,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  goto do_notice;
 
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
 	  goto do_add;
@@ -7327,7 +7350,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  break;
 
 	case OMP_CLAUSE_DEVICE_RESIDENT:
-	case OMP_CLAUSE_USE_DEVICE:
 	  remove = true;
 	  break;
 
@@ -9365,6 +9387,9 @@  gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     case OMP_TEAMS:
       ort = OMP_TEAMS_COMBINED (expr) ? ORT_COMBINED_TEAMS : ORT_TEAMS;
       break;
+    case OACC_HOST_DATA:
+      ort = ORT_ACC_HOST_DATA;
+      break;
     default:
       gcc_unreachable ();
     }
@@ -9386,6 +9411,7 @@  gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 	  switch (TREE_CODE (expr))
 	    {
 	    case OACC_DATA:
+	    case OACC_HOST_DATA:
 	      end_ix = BUILT_IN_GOACC_DATA_END;
 	      break;
 	    case OMP_TARGET_DATA:
@@ -9418,6 +9444,10 @@  gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
       stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
 				      OMP_CLAUSES (expr));
       break;
+    case OACC_HOST_DATA:
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA,
+				      OMP_CLAUSES (expr));
+      break;
     case OACC_PARALLEL:
       stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
 				      OMP_CLAUSES (expr));
@@ -10527,16 +10557,12 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = GS_ALL_DONE;
 	  break;
 
-	case OACC_HOST_DATA:
-	  sorry ("directive not yet implemented");
-	  ret = GS_ALL_DONE;
-	  break;
-
 	case OACC_DECLARE:
 	  gimplify_oacc_declare (expr_p, pre_p);
 	  ret = GS_ALL_DONE;
 	  break;
 
+	case OACC_HOST_DATA:
 	case OACC_DATA:
 	case OACC_KERNELS:
 	case OACC_PARALLEL:
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index d540dab..35f5014 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -47,6 +47,8 @@  DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait",
 		   BT_FN_VOID_INT_INT_VAR,
 		   ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_HOST_DATA, "GOACC_host_data",
+		   BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
 
 DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
 			    BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index f17a828..15cc839 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1942,6 +1942,7 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  install_var_local (decl, ctx);
 	  break;
 
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
@@ -2144,7 +2145,6 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE_DEVICE_RESIDENT:
-	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE__CACHE_:
 	  sorry ("Clause not supported yet");
 	  break;
@@ -2295,6 +2295,7 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_NOGROUP:
 	case OMP_CLAUSE_DEFAULTMAP:
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	case OMP_CLAUSE_ASYNC:
@@ -2312,7 +2313,6 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE_DEVICE_RESIDENT:
-	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE__CACHE_:
 	  sorry ("Clause not supported yet");
 	  break;
@@ -3615,6 +3615,8 @@  check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 	    case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
 	    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
 	      stmt_name = "enter/exit data"; break;
+	    case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
+	      break;
 	    default: gcc_unreachable ();
 	    }
 	  switch (gimple_omp_target_kind (ctx->stmt))
@@ -3626,6 +3628,8 @@  check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 	    case GF_OMP_TARGET_KIND_OACC_KERNELS:
 	      ctx_stmt_name = "kernels"; break;
 	    case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
+	    case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+	      ctx_stmt_name = "host_data"; break;
 	    default: gcc_unreachable ();
 	    }
 
@@ -12508,6 +12512,7 @@  expand_omp_target (struct omp_region *region)
       break;
     case GF_OMP_TARGET_KIND_DATA:
     case GF_OMP_TARGET_KIND_OACC_DATA:
+    case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
       data_region = true;
       break;
     default:
@@ -12751,6 +12756,9 @@  expand_omp_target (struct omp_region *region)
     case GF_OMP_TARGET_KIND_OACC_DECLARE:
       start_ix = BUILT_IN_GOACC_DECLARE;
       break;
+    case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+      start_ix = BUILT_IN_GOACC_HOST_DATA;
+      break;
     default:
       gcc_unreachable ();
     }
@@ -12875,6 +12883,7 @@  expand_omp_target (struct omp_region *region)
     case BUILT_IN_GOACC_DATA_START:
     case BUILT_IN_GOACC_DECLARE:
     case BUILT_IN_GOMP_TARGET_DATA:
+    case BUILT_IN_GOACC_HOST_DATA:
       break;
     case BUILT_IN_GOMP_TARGET:
     case BUILT_IN_GOMP_TARGET_UPDATE:
@@ -13182,6 +13191,7 @@  build_omp_regions_1 (basic_block bb, struct omp_region *parent,
 		case GF_OMP_TARGET_KIND_OACC_PARALLEL:
 		case GF_OMP_TARGET_KIND_OACC_KERNELS:
 		case GF_OMP_TARGET_KIND_OACC_DATA:
+		case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
 		  break;
 		case GF_OMP_TARGET_KIND_UPDATE:
 		case GF_OMP_TARGET_KIND_ENTER_DATA:
@@ -14982,6 +14992,7 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       break;
     case GF_OMP_TARGET_KIND_DATA:
     case GF_OMP_TARGET_KIND_OACC_DATA:
+    case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
       data_region = true;
       break;
     default:
@@ -15188,6 +15199,7 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  }
 	break;
 
+      case OMP_CLAUSE_USE_DEVICE:
       case OMP_CLAUSE_USE_DEVICE_PTR:
       case OMP_CLAUSE_IS_DEVICE_PTR:
 	var = OMP_CLAUSE_DECL (c);
@@ -15573,12 +15585,14 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				    build_int_cstu (tkind_type, tkind));
 	    break;
 
+	  case OMP_CLAUSE_USE_DEVICE:
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    ovar = OMP_CLAUSE_DECL (c);
 	    var = lookup_decl_in_outer_ctx (ovar, ctx);
 	    x = build_sender_ref (ovar, ctx);
-	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+		|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
 	      tkind = GOMP_MAP_USE_DEVICE_PTR;
 	    else
 	      tkind = GOMP_MAP_FIRSTPRIVATE_INT;
@@ -15781,10 +15795,12 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				     gimple_build_assign (new_var, x));
 	      }
 	    break;
+	  case OMP_CLAUSE_USE_DEVICE:
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    var = OMP_CLAUSE_DECL (c);
-	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+		|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
 	      x = build_sender_ref (var, ctx);
 	    else
 	      x = build_receiver_ref (var, false, ctx);
@@ -16771,6 +16787,7 @@  make_gimple_omp_edges (basic_block bb, struct omp_region **region,
 	case GF_OMP_TARGET_KIND_OACC_PARALLEL:
 	case GF_OMP_TARGET_KIND_OACC_KERNELS:
 	case GF_OMP_TARGET_KIND_OACC_DATA:
+	case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
 	  break;
 	case GF_OMP_TARGET_KIND_UPDATE:
 	case GF_OMP_TARGET_KIND_ENTER_DATA:
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 280d29b..70904ce 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1072,6 +1072,7 @@  convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SHARED:
 	case OMP_CLAUSE_TO_DECLARE:
 	case OMP_CLAUSE_LINK:
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	do_decl_clause:
@@ -1721,6 +1722,7 @@  convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SHARED:
 	case OMP_CLAUSE_TO_DECLARE:
 	case OMP_CLAUSE_LINK:
+	case OMP_CLAUSE_USE_DEVICE:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	do_decl_clause:
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 4d42c42..ea9344d 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -394,6 +394,7 @@  GOACC_2.0.1 {
   global:
 	GOACC_declare;
 	GOACC_parallel_keyed;
+	GOACC_host_data;
 } GOACC_2.0;
 
 GOMP_PLUGIN_1.0 {
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index a80ede4..db7cab3 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -490,6 +490,46 @@  GOACC_wait (int async, int num_waits, ...)
     goacc_thread ()->dev->openacc.async_wait_all_async_func (acc_async_noval);
 }
 
+void
+GOACC_host_data (int device, size_t mapnum,
+		 void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+  bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
+  struct target_mem_desc *tgt;
+
+#ifdef HAVE_INTTYPES_H
+  gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n",
+	      __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds);
+#else
+  gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p\n",
+	      __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds);
+#endif
+
+  goacc_lazy_initialize ();
+
+  struct goacc_thread *thr = goacc_thread ();
+  struct gomp_device_descr *acc_dev = thr->dev;
+
+  /* Host fallback or 'do nothing'.  */
+  if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+      || host_fallback)
+    {
+      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
+			   GOMP_MAP_VARS_OPENACC);
+      tgt->prev = thr->mapped_data;
+      thr->mapped_data = tgt;
+
+      return;
+    }
+
+  gomp_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
+  tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
+		       GOMP_MAP_VARS_OPENACC);
+  gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
+  tgt->prev = thr->mapped_data;
+  thr->mapped_data = tgt;
+}
+
 int
 GOACC_get_num_threads (void)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
new file mode 100644
index 0000000..51745ba
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
@@ -0,0 +1,100 @@ 
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-lcuda -lcublas -lcudart" } */
+
+#include <stdlib.h>
+#include <openacc.h>
+#include <cuda.h>
+#include <cuda_runtime_api.h>
+#include <cublas_v2.h>
+
+void
+saxpy_host (int n, float a, float *x, float *y)
+{
+  int i;
+
+  for (i = 0; i < n; i++)
+    y[i] = y[i] + a * x[i];
+}
+
+#pragma acc routine
+void
+saxpy_target (int n, float a, float *x, float *y)
+{
+  int i;
+
+  for (i = 0; i < n; i++)
+    y[i] = y[i] + a * x[i];
+}
+
+int
+main(int argc, char **argv)
+{
+#define N 8
+  int i;
+  float x_ref[N], y_ref[N];
+  float x[N], y[N];
+  cublasHandle_t h;
+  float a = 2.0;
+
+  for (i = 0; i < N; i++)
+    {
+      x[i] = x_ref[i] = 4.0 + i;
+      y[i] = y_ref[i] = 3.0;
+    }
+
+  saxpy_host (N, a, x_ref, y_ref);
+
+  cublasCreate (&h);
+
+#pragma acc data copyin (x[0:N]) copy (y[0:N])
+  {
+#pragma acc host_data use_device (x, y)
+    {
+      cublasSaxpy (h, N, &a, x, 1, y, 1);
+    }
+  }
+
+  for (i = 0; i < N; i++)
+    {
+      if (y[i] != y_ref[i])
+        abort ();
+    }
+
+#pragma acc data create (x[0:N]) copyout (y[0:N])
+  {
+#pragma acc kernels
+    for (i = 0; i < N; i++)
+      y[i] = 3.0;
+
+#pragma acc host_data use_device (x, y)
+    {
+      cublasSaxpy (h, N, &a, x, 1, y, 1);
+    }
+  }
+
+  cublasDestroy (h);
+
+  for (i = 0; i < N; i++)
+    {
+      if (y[i] != y_ref[i])
+        abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    y[i] = 3.0;
+
+  /* There's no need to use host_data here.  */
+#pragma acc data copyin (x[0:N]) copyin (a) copy (y[0:N])
+  {
+#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a)
+    saxpy_target (N, a, x, y);
+  }
+
+  for (i = 0; i < N; i++)
+    {
+      if (y[i] != y_ref[i])
+        abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
new file mode 100644
index 0000000..9820286
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
@@ -0,0 +1,31 @@ 
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+char *global_in_host;
+
+void foo (char *in)
+{
+  if (!acc_is_present (global_in_host, sizeof (*global_in_host))
+      || in != acc_deviceptr (global_in_host))
+    abort ();
+}
+
+int
+main (int argc, char **argv)
+{
+  char mydata[1024];
+
+  global_in_host = mydata;
+
+#pragma acc data copyin(mydata)
+  {
+#pragma acc host_data use_device (mydata)
+    {
+      foo (mydata);
+    }
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c
new file mode 100644
index 0000000..b6ee9b1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c
@@ -0,0 +1,29 @@ 
+/* { dg-do compile } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+  int x[N];
+
+#pragma acc data copyin (x[0:N])
+  {
+    int *xp;
+#pragma acc host_data use_device (x)
+    {
+      /* This use of the present clause is undefined behaviour for OpenACC.  */
+#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing host_data region" } */
+      {
+        xp = x;
+      }
+    }
+
+    if (xp != acc_deviceptr (x))
+      abort ();
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c
new file mode 100644
index 0000000..3504f27
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c
@@ -0,0 +1,29 @@ 
+/* { dg-do run } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+  int x[N], *xp2;
+
+#pragma acc data copyin (x[0:N])
+  {
+    int *xp;
+#pragma acc host_data use_device (x)
+    {
+#pragma acc data
+      {
+        xp = x;
+      }
+      xp2 = x;
+    }
+
+    if (xp != acc_deviceptr (x) || xp2 != xp)
+      abort ();
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c
new file mode 100644
index 0000000..268e919
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c
@@ -0,0 +1,38 @@ 
+/* { dg-do run } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+  int x[N], y[N], *yp;
+
+  yp = y + 1;
+
+#pragma acc data copyin (x[0:N])
+  {
+    int *xp, *yp2;
+#pragma acc host_data use_device (x)
+    {
+#pragma acc data copyin (y)
+      {
+#pragma acc host_data use_device (yp)
+	{
+	  xp = x;
+	  yp2 = yp;
+	}
+
+        if (yp2 != acc_deviceptr (yp))
+	  abort ();
+      }
+    }
+
+    if (xp != acc_deviceptr (x))
+      abort ();
+
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
new file mode 100644
index 0000000..d0b1968
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
@@ -0,0 +1,31 @@ 
+/* { dg-do compile } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+  int x[N];
+
+#pragma acc data copyin (x[0:N])
+  {
+    int *xp;
+#pragma acc host_data use_device (x)
+    {
+      /* Here 'x' being implicitly firstprivate for the parallel region
+	 conflicts with it being declared as use_device in the enclosing
+	 host_data region.  */
+#pragma acc parallel copyout (xp)
+      {
+        xp = x; /* { dg-error "variable 'x' declared in enclosing host_data region" } */
+      }
+    }
+
+    if (xp != acc_deviceptr (x))
+      abort ();
+  }
+
+  return 0;
+}