[og8] Add OpenACC 2.6 if and if_present clauses on host_data construct

OpenACC 2.6 specifies `if' and `if_present' clauses on the `host_data' 
construct. These patches add support for these clauses. The first patch, 
by Thomas, reorganizes libgomp internals to turn a "device" argument 
into "flags" that can provide more information to the runtime. The 
second patch adds support for the `if' and `if_present' clauses, using 
the new flag mechanism.

OK for openacc-gcc-8-branch?

     gcc/
     * omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP
     code paths.  Update for libgomp OpenACC entry points change.
     include/
     * gomp-constants.h (GOACC_FLAG_HOST_FALLBACK)
     (GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define.
     libgomp/
     * oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel)
     (GOACC_data_start, GOACC_enter_exit_data, GOACC_update)
     (GOACC_declare): Redefine the "device" argument to "flags".


     gcc/c/
     * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
     and PRAGMA_OACC_CLAUSE_IF_PRESENT.
     gcc/cp/
     * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Likewise.

     gcc/fortran/
     * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
     OMP_CLAUSE_IF_PRESENT.

     gcc/
     * omp-expand.c (expand_omp_target): Handle if_present flag on
     OpenACC host_data construct.

     gcc/testsuite/c-c++-common/goacc/
     * host_data-1.c: Add tests of if and if_present clauses on host_data.
     gcc/testsuite/gfortran.dg/goacc/
     * host_data-tree.f95: Likewise.

     include/
     * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant.

     libgomp/
     * libgomp.h (enum gomp_map_vars_kind): Add
     GOMP_MAP_VARS_OPENACC_IF_PRESENT.

     libgomp/
     * oacc-parallel.c (GOACC_data_start): Handle
     GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
     * target.c (gomp_map_vars_async): Handle
     GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.

     libgomp/testsuite/libgomp.oacc-c-c++-common/
     * host_data-6.c: New test.

Hi Gergő!

On Fri, 21 Dec 2018 13:29:09 +0100, Gergö Barany <gergo@codesourcery.com> wrote:
> OpenACC 2.6 specifies `if' and `if_present' clauses on the `host_data' 
> construct. These patches add support for these clauses. The first patch, 
> by Thomas, reorganizes libgomp internals to turn a "device" argument 
> into "flags" that can provide more information to the runtime. The 
> second patch adds support for the `if' and `if_present' clauses, using 
> the new flag mechanism.
> 
> OK for openacc-gcc-8-branch?

Yes, thanks.  To record the review effort, please include "Reviewed-by:
Thomas Schwinge <thomas@codesourcery.com>" in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.  (Not for my own commit, of
course.)

Again, just the commit message of the second commit needs to be adjusted, from:

>     [...]
>     gcc/testsuite/c-c++-common/goacc/
>     * host_data-1.c: Add tests of if and if_present clauses on host_data.
>     gcc/testsuite/gfortran.dg/goacc/
>     * host_data-tree.f95: Likewise.
>     [...]
>     libgomp/
>     * libgomp.h (enum gomp_map_vars_kind): Add
>     GOMP_MAP_VARS_OPENACC_IF_PRESENT.
> 
>     libgomp/
>     * oacc-parallel.c (GOACC_data_start): Handle
>     GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
>     * target.c (gomp_map_vars_async): Handle
>     GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.
> 
>     libgomp/testsuite/libgomp.oacc-c-c++-common/
>     * host_data-6.c: New test.

... to:

>     [...]
>     gcc/testsuite/
>     * c-c++-common/goacc/host_data-1.c: Add tests of if and if_present clauses on host_data. [add suitable line break some where]
>     * gfortran.dg/goacc/host_data-tree.f95: Likewise.
>     [...]
>     libgomp/
>     * libgomp.h (enum gomp_map_vars_kind): Add
>     GOMP_MAP_VARS_OPENACC_IF_PRESENT.
>     * oacc-parallel.c (GOACC_data_start): Handle
>     GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
>     * target.c (gomp_map_vars_async): Handle
>     GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.
>     * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.


Grüße
 Thomas


> From 6d719cc2bcfa8f7ed8cb59e753e44aab6bf634fb Mon Sep 17 00:00:00 2001
> From: Thomas Schwinge <thomas@codesourcery.com>
> Date: Wed, 19 Dec 2018 20:04:18 +0100
> Subject: [PATCH 1/2] For libgomp OpenACC entry points, redefine the "device"
>  argument to "flags"
> 
> ... so that we're then able to use this for other flags in addition to
> "GOACC_FLAG_HOST_FALLBACK".
> 
> 	gcc/
> 	* omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP
> 	code paths.  Update for libgomp OpenACC entry points change.
> 	include/
> 	* gomp-constants.h (GOACC_FLAG_HOST_FALLBACK)
> 	(GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define.
> 	libgomp/
> 	* oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel)
> 	(GOACC_data_start, GOACC_enter_exit_data, GOACC_update)
> 	(GOACC_declare): Redefine the "device" argument to "flags".
> ---
>  gcc/ChangeLog.openacc      |   5 ++
>  gcc/omp-expand.c           | 111 +++++++++++++++++++++++++++++----------------
>  gcc/tree-ssa-structalias.c |   4 +-
>  include/ChangeLog.openacc  |   5 ++
>  include/gomp-constants.h   |  12 +++++
>  libgomp/ChangeLog.openacc  |   6 +++
>  libgomp/oacc-parallel.c    |  60 ++++++++++++++----------
>  7 files changed, 139 insertions(+), 64 deletions(-)
> 
> diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
> index 718044c..6a51b1e 100644
> --- a/gcc/ChangeLog.openacc
> +++ b/gcc/ChangeLog.openacc
> @@ -1,3 +1,8 @@
> +2018-12-21  Thomas Schwinge  <thomas@codesourcery.com>
> +
> +	* omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP
> +	code paths.  Update for libgomp OpenACC entry points change.
> +
>  2018-12-21  Gergö Barany  <gergo@codesourcery.com>
>  
>  	* omp-low.c (scan_sharing_clauses): Fix call to renamed function
> diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
> index 988b1bb..ea264da 100644
> --- a/gcc/omp-expand.c
> +++ b/gcc/omp-expand.c
> @@ -7204,7 +7204,7 @@ expand_omp_target (struct omp_region *region)
>       transfers.  */
>    tree t1, t2, t3, t4, device, cond, depend, c, clauses;
>    enum built_in_function start_ix;
> -  location_t clause_loc;
> +  location_t clause_loc = UNKNOWN_LOCATION;
>    unsigned int flags_i = 0;
>  
>    switch (gimple_omp_target_kind (entry_stmt))
> @@ -7249,49 +7249,62 @@ expand_omp_target (struct omp_region *region)
>  
>    clauses = gimple_omp_target_clauses (entry_stmt);
>  
> -  /* By default, the value of DEVICE is GOMP_DEVICE_ICV (let runtime
> -     library choose) and there is no conditional.  */
> -  cond = NULL_TREE;
> -  device = build_int_cst (integer_type_node, GOMP_DEVICE_ICV);
> -
> -  c = omp_find_clause (clauses, OMP_CLAUSE_IF);
> -  if (c)
> -    cond = OMP_CLAUSE_IF_EXPR (c);
> -
> -  c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE);
> -  if (c)
> +  device = NULL_TREE;
> +  tree goacc_flags = NULL_TREE;
> +  if (is_gimple_omp_oacc (entry_stmt))
>      {
> -      /* Even if we pass it to all library function calls, it is currently only
> -	 defined/used for the OpenMP target ones.  */
> -      gcc_checking_assert (start_ix == BUILT_IN_GOMP_TARGET
> -			   || start_ix == BUILT_IN_GOMP_TARGET_DATA
> -			   || start_ix == BUILT_IN_GOMP_TARGET_UPDATE
> -			   || start_ix == BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA);
> -
> -      device = OMP_CLAUSE_DEVICE_ID (c);
> -      clause_loc = OMP_CLAUSE_LOCATION (c);
> +      /* By default, no GOACC_FLAGs are set.  */
> +      goacc_flags = integer_zero_node;
>      }
>    else
> -    clause_loc = gimple_location (entry_stmt);
> -
> -  c = omp_find_clause (clauses, OMP_CLAUSE_NOWAIT);
> -  if (c)
> -    flags_i |= GOMP_TARGET_FLAG_NOWAIT;
> +    {
> +      c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE);
> +      if (c)
> +	{
> +	  device = OMP_CLAUSE_DEVICE_ID (c);
> +	  clause_loc = OMP_CLAUSE_LOCATION (c);
> +	}
> +      else
> +	{
> +	  /* By default, the value of DEVICE is GOMP_DEVICE_ICV (let runtime
> +	     library choose).  */
> +	  device = build_int_cst (integer_type_node, GOMP_DEVICE_ICV);
> +	  clause_loc = gimple_location (entry_stmt);
> +	}
>  
> -  /* Ensure 'device' is of the correct type.  */
> -  device = fold_convert_loc (clause_loc, integer_type_node, device);
> +      c = omp_find_clause (clauses, OMP_CLAUSE_NOWAIT);
> +      if (c)
> +	flags_i |= GOMP_TARGET_FLAG_NOWAIT;
> +    }
>  
> -  /* If we found the clause 'if (cond)', build
> -     (cond ? device : GOMP_DEVICE_HOST_FALLBACK).  */
> +  /* By default, there is no conditional.  */
> +  cond = NULL_TREE;
> +  c = omp_find_clause (clauses, OMP_CLAUSE_IF);
> +  if (c)
> +    cond = OMP_CLAUSE_IF_EXPR (c);
> +  /* If we found the clause 'if (cond)', build:
> +     OpenACC: goacc_flags = (cond ? goacc_flags : flags | GOACC_FLAG_HOST_FALLBACK)
> +     OpenMP: device = (cond ? device : GOMP_DEVICE_HOST_FALLBACK) */
>    if (cond)
>      {
> +      tree *tp;
> +      if (is_gimple_omp_oacc (entry_stmt))
> +	tp = &goacc_flags;
> +      else
> +	{
> +	  /* Ensure 'device' is of the correct type.  */
> +	  device = fold_convert_loc (clause_loc, integer_type_node, device);
> +
> +	  tp = &device;
> +	}
> +
>        cond = gimple_boolify (cond);
>  
>        basic_block cond_bb, then_bb, else_bb;
>        edge e;
>        tree tmp_var;
>  
> -      tmp_var = create_tmp_var (TREE_TYPE (device));
> +      tmp_var = create_tmp_var (TREE_TYPE (*tp));
>        if (offloaded)
>  	e = split_block_after_labels (new_bb);
>        else
> @@ -7314,13 +7327,20 @@ expand_omp_target (struct omp_region *region)
>        gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
>  
>        gsi = gsi_start_bb (then_bb);
> -      stmt = gimple_build_assign (tmp_var, device);
> +      stmt = gimple_build_assign (tmp_var, *tp);
>        gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
>  
>        gsi = gsi_start_bb (else_bb);
> -      stmt = gimple_build_assign (tmp_var,
> -				  build_int_cst (integer_type_node,
> -						 GOMP_DEVICE_HOST_FALLBACK));
> +      if (is_gimple_omp_oacc (entry_stmt))
> +	stmt = gimple_build_assign (tmp_var,
> +				    BIT_IOR_EXPR,
> +				    *tp,
> +				    build_int_cst (integer_type_node,
> +						   GOACC_FLAG_HOST_FALLBACK));
> +      else
> +	stmt = gimple_build_assign (tmp_var,
> +				    build_int_cst (integer_type_node,
> +						   GOMP_DEVICE_HOST_FALLBACK));
>        gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
>  
>        make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
> @@ -7330,14 +7350,17 @@ expand_omp_target (struct omp_region *region)
>        make_edge (then_bb, new_bb, EDGE_FALLTHRU);
>        make_edge (else_bb, new_bb, EDGE_FALLTHRU);
>  
> -      device = tmp_var;
> +      *tp = tmp_var;
> +
>        gsi = gsi_last_nondebug_bb (new_bb);
>      }
>    else
>      {
>        gsi = gsi_last_nondebug_bb (new_bb);
> -      device = force_gimple_operand_gsi (&gsi, device, true, NULL_TREE,
> -					 true, GSI_SAME_STMT);
> +
> +      if (device != NULL_TREE)
> +	device = force_gimple_operand_gsi (&gsi, device, true, NULL_TREE,
> +					   true, GSI_SAME_STMT);
>      }
>  
>    t = gimple_omp_target_data_arg (entry_stmt);
> @@ -7361,7 +7384,17 @@ expand_omp_target (struct omp_region *region)
>    bool tagging = false;
>    /* The maximum number used by any start_ix, without varargs.  */
>    auto_vec<tree, 11> args;
> -  args.quick_push (device);
> +  if (is_gimple_omp_oacc (entry_stmt))
> +    {
> +      tree goacc_flags_m = fold_build1 (GOACC_FLAGS_MARSHAL_OP,
> +					TREE_TYPE (goacc_flags), goacc_flags);
> +      goacc_flags_m = force_gimple_operand_gsi (&gsi, goacc_flags_m, true,
> +						NULL_TREE, true,
> +						GSI_SAME_STMT);
> +      args.quick_push (goacc_flags_m);
> +    }
> +  else
> +    args.quick_push (device);
>    if (start_ix == BUILT_IN_GOACC_PARALLEL)
>      {
>        tree use_params = oacc_parallel ? integer_one_node : integer_zero_node;
> diff --git a/gcc/tree-ssa-structalias.c b/gcc/tree-ssa-structalias.c
> index a4f7251..bcf3fd3 100644
> --- a/gcc/tree-ssa-structalias.c
> +++ b/gcc/tree-ssa-structalias.c
> @@ -4684,7 +4684,7 @@ find_func_aliases_for_builtin_call (struct function *fn, gcall *t)
>  		  argpos = 1;
>  		  break;
>  		case BUILT_IN_GOACC_PARALLEL:
> -		  /* __builtin_GOACC_parallel (device, fn, mapnum, hostaddrs,
> +		  /* __builtin_GOACC_parallel (flags_m, fn, mapnum, hostaddrs,
>  					       sizes, kinds, ...).  */
>  		  fnpos = 2;
>  		  argpos = 4;
> @@ -5263,7 +5263,7 @@ find_func_clobbers (struct function *fn, gimple *origt)
>  		  argpos = 1;
>  		  break;
>  		case BUILT_IN_GOACC_PARALLEL:
> -		  /* __builtin_GOACC_parallel (device, fn, mapnum, hostaddrs,
> +		  /* __builtin_GOACC_parallel (flags_m, fn, mapnum, hostaddrs,
>  					       sizes, kinds, ...).  */
>  		  fnpos = 2;
>  		  argpos = 4;
> diff --git a/include/ChangeLog.openacc b/include/ChangeLog.openacc
> index 20ed27f..aa583ea 100644
> --- a/include/ChangeLog.openacc
> +++ b/include/ChangeLog.openacc
> @@ -1,3 +1,8 @@
> +2018-12-21  Thomas Schwinge  <thomas@codesourcery.com>
> +
> +	* gomp-constants.h (GOACC_FLAG_HOST_FALLBACK)
> +	(GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define.
> +
>  2018-12-20  Julian Brown  <julian@codesourcery.com>
>  	    Maciej W. Rozycki  <macro@codesourcery.com>
>  
> diff --git a/include/gomp-constants.h b/include/gomp-constants.h
> index 27de5bc..b5d8441 100644
> --- a/include/gomp-constants.h
> +++ b/include/gomp-constants.h
> @@ -252,6 +252,18 @@ enum gomp_map_kind
>  /* Internal to libgomp.  */
>  #define GOMP_TARGET_FLAG_UPDATE		(1U << 31)
>  
> +
> +/* OpenACC construct flags.  */
> +
> +/* Force host fallback execution.  */
> +#define GOACC_FLAG_HOST_FALLBACK	(1 << 0)
> +
> +/* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted
> +   bitmask.  */
> +#define GOACC_FLAGS_MARSHAL_OP		BIT_NOT_EXPR
> +#define GOACC_FLAGS_UNMARSHAL(X)	(~(X))
> +
> +
>  /* Versions of libgomp and device-specific plugins.  GOMP_VERSION
>     should be incremented whenever an ABI-incompatible change is introduced
>     to the plugin interface defined in libgomp/libgomp.h.  */
> diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc
> index b48453b..04cea5f 100644
> --- a/libgomp/ChangeLog.openacc
> +++ b/libgomp/ChangeLog.openacc
> @@ -1,3 +1,9 @@
> +2018-12-21  Thomas Schwinge  <thomas@codesourcery.com>
> +
> +	* oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel)
> +	(GOACC_data_start, GOACC_enter_exit_data, GOACC_update)
> +	(GOACC_declare): Redefine the "device" argument to "flags".
> +
>  2018-12-20  Gergö Barany  <gergo@codesourcery.com>
>  	    Thomas Schwinge  <thomas@codesourcery.com>
>  
> diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
> index c74221f..0b5f41a 100644
> --- a/libgomp/oacc-parallel.c
> +++ b/libgomp/oacc-parallel.c
> @@ -41,6 +41,16 @@
>  #include <stdarg.h>
>  #include <assert.h>
>  
> +
> +/* In the ABI, the GOACC_FLAGs are encoded as an inverted bitmask, so that we
> +   continue to support the following two legacy values.  */
> +_Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_ICV) == 0,
> +		"legacy GOMP_DEVICE_ICV broken");
> +_Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_HOST_FALLBACK)
> +		== GOACC_FLAG_HOST_FALLBACK,
> +		"legacy GOMP_DEVICE_HOST_FALLBACK broken");
> +
> +
>  /* Returns the number of mappings associated with the pointer or pset. PSET
>     have three mappings, whereas pointer have two.  */
>  
> @@ -159,17 +169,18 @@ goacc_call_host_fn (void (*fn) (void *), size_t mapnum, void **hostaddrs,
>    fn (hostaddrs);
>  }
>  
> -/* Launch a possibly offloaded function on DEVICE.  FN is the host fn
> +/* Launch a possibly offloaded function with FLAGS.  FN is the host fn
>     address.  MAPNUM, HOSTADDRS, SIZES & KINDS  describe the memory
>     blocks to be copied to/from the device.  Varadic arguments are
>     keyed optional parameters terminated with a zero.  */
>  
>  static void
> -GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
> +GOACC_parallel_keyed_internal (int flags_m, int params, void (*fn) (void *),
>  			       size_t mapnum, void **hostaddrs, size_t *sizes,
>  			       unsigned short *kinds, va_list *ap)
>  {
> -  bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
> +  int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
> +
>    struct goacc_thread *thr;
>    struct gomp_device_descr *acc_dev;
>    struct target_mem_desc *tgt;
> @@ -252,7 +263,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
>  
>    /* Host fallback if "if" clause is false or if the current device is set to
>       the host.  */
> -  if (host_fallback)
> +  if (flags & GOACC_FLAG_HOST_FALLBACK)
>      {
>        //TODO
>        prof_info.device_type = acc_device_host;
> @@ -448,25 +459,25 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
>  }
>  
>  void
> -GOACC_parallel_keyed (int device, void (*fn) (void *),
> +GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
>  		      size_t mapnum, void **hostaddrs, size_t *sizes,
>  		      unsigned short *kinds, ...)
>  {
>    va_list ap;
>    va_start (ap, kinds);
> -  GOACC_parallel_keyed_internal (device, 0, fn, mapnum, hostaddrs, sizes,
> +  GOACC_parallel_keyed_internal (flags_m, 0, fn, mapnum, hostaddrs, sizes,
>  				 kinds, &ap);
>    va_end (ap);
>  }
>  
>  void
> -GOACC_parallel_keyed_v2 (int device, int args, void (*fn) (void *),
> +GOACC_parallel_keyed_v2 (int flags_m, int args, void (*fn) (void *),
>  			 size_t mapnum, void **hostaddrs, size_t *sizes,
>  			 unsigned short *kinds, ...)
>  {
>    va_list ap;
>    va_start (ap, kinds);
> -  GOACC_parallel_keyed_internal (device, args, fn, mapnum, hostaddrs, sizes,
> +  GOACC_parallel_keyed_internal (flags_m, args, fn, mapnum, hostaddrs, sizes,
>  				 kinds, &ap);
>    va_end (ap);
>  }
> @@ -474,7 +485,7 @@ GOACC_parallel_keyed_v2 (int device, int args, void (*fn) (void *),
>  /* Legacy entry point, only provide host execution.  */
>  
>  void
> -GOACC_parallel (int device, void (*fn) (void *),
> +GOACC_parallel (int flags_m, void (*fn) (void *),
>  		size_t mapnum, void **hostaddrs, size_t *sizes,
>  		unsigned short *kinds,
>  		int num_gangs, int num_workers, int vector_length,
> @@ -486,10 +497,11 @@ GOACC_parallel (int device, void (*fn) (void *),
>  }
>  
>  void
> -GOACC_data_start (int device, size_t mapnum,
> +GOACC_data_start (int flags_m, size_t mapnum,
>  		  void **hostaddrs, size_t *sizes, unsigned short *kinds)
>  {
> -  bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
> +  int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
> +
>    struct target_mem_desc *tgt;
>  
>  #ifdef HAVE_INTTYPES_H
> @@ -575,7 +587,7 @@ GOACC_data_start (int device, size_t mapnum,
>  
>    /* Host fallback or 'do nothing'.  */
>    if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
> -      || host_fallback)
> +      || (flags & GOACC_FLAG_HOST_FALLBACK))
>      {
>        //TODO
>        prof_info.device_type = acc_device_host;
> @@ -694,13 +706,14 @@ GOACC_data_end (void)
>  }
>  
>  void
> -GOACC_enter_exit_data (int device, size_t mapnum,
> +GOACC_enter_exit_data (int flags_m, size_t mapnum,
>  		       void **hostaddrs, size_t *sizes, unsigned short *kinds,
>  		       int async, int num_waits, ...)
>  {
> +  int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
> +
>    struct goacc_thread *thr;
>    struct gomp_device_descr *acc_dev;
> -  bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
>    bool data_enter = false;
>    size_t i;
>  
> @@ -815,7 +828,7 @@ GOACC_enter_exit_data (int device, size_t mapnum,
>  			      &api_info);
>  
>    if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
> -      || host_fallback)
> +      || (flags & GOACC_FLAG_HOST_FALLBACK))
>      {
>        //TODO
>        prof_info.device_type = acc_device_host;
> @@ -1098,11 +1111,12 @@ goacc_wait (int async, int num_waits, va_list *ap)
>  }
>  
>  void
> -GOACC_update (int device, size_t mapnum,
> +GOACC_update (int flags_m, size_t mapnum,
>  	      void **hostaddrs, size_t *sizes, unsigned short *kinds,
>  	      int async, int num_waits, ...)
>  {
> -  bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
> +  int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
> +
>    size_t i;
>  
>    goacc_lazy_initialize ();
> @@ -1163,7 +1177,7 @@ GOACC_update (int device, size_t mapnum,
>      goacc_profiling_dispatch (&prof_info, &update_event_info, &api_info);
>  
>    if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
> -      || host_fallback)
> +      || (flags & GOACC_FLAG_HOST_FALLBACK))
>      {
>        //TODO
>        prof_info.device_type = acc_device_host;
> @@ -1309,7 +1323,7 @@ GOACC_get_thread_num (void)
>  }
>  
>  void
> -GOACC_declare (int device, size_t mapnum,
> +GOACC_declare (int flags_m, size_t mapnum,
>  	       void **hostaddrs, size_t *sizes, unsigned short *kinds)
>  {
>    int i;
> @@ -1329,7 +1343,7 @@ GOACC_declare (int device, size_t mapnum,
>  	  case GOMP_MAP_POINTER:
>  	  case GOMP_MAP_RELEASE:
>  	  case GOMP_MAP_DELETE:
> -	    GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
> +	    GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i],
>  				   &kinds[i], 0, 0);
>  	    break;
>  
> @@ -1338,18 +1352,18 @@ GOACC_declare (int device, size_t mapnum,
>  
>  	  case GOMP_MAP_ALLOC:
>  	    if (!acc_is_present (hostaddrs[i], sizes[i]))
> -	      GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
> +	      GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i],
>  				     &kinds[i], 0, 0);
>  	    break;
>  
>  	  case GOMP_MAP_TO:
> -	    GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
> +	    GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i],
>  				   &kinds[i], 0, 0);
>  
>  	    break;
>  
>  	  case GOMP_MAP_FROM:
> -	    GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
> +	    GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i],
>  				   &kinds[i], 0, 0);
>  	    break;
>  
> -- 
> 2.8.1
> 
> From cbd9efcd4ebb6c73a14ead01d85e452d63b7c937 Mon Sep 17 00:00:00 2001
> From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@codesourcery.com>
> Date: Fri, 21 Dec 2018 01:12:44 -0800
> Subject: [PATCH 2/2] [og8] Add OpenACC 2.6 if and if_present clauses on
>  host_data construct: GOACC_FLAG_HOST_DATA_IF_PRESENT
> 
>     gcc/c/
>     * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
>     and PRAGMA_OACC_CLAUSE_IF_PRESENT.
>     gcc/cp/
>     * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Likewise.
> 
>     gcc/fortran/
>     * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
>     OMP_CLAUSE_IF_PRESENT.
> 
>     gcc/
>     * omp-expand.c (expand_omp_target): Handle if_present flag on
>     OpenACC host_data construct.
> 
>     gcc/testsuite/c-c++-common/goacc/
>     * host_data-1.c: Add tests of if and if_present clauses on host_data.
>     gcc/testsuite/gfortran.dg/goacc/
>     * host_data-tree.f95: Likewise.
> 
>     include/
>     * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant.
> 
>     libgomp/
>     * libgomp.h (enum gomp_map_vars_kind): Add
>     GOMP_MAP_VARS_OPENACC_IF_PRESENT.
> 
>     libgomp/
>     * oacc-parallel.c (GOACC_data_start): Handle
>     GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
>     * target.c (gomp_map_vars_async): Handle
>     GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.
> 
>     libgomp/testsuite/libgomp.oacc-c-c++-common/
>     * host_data-6.c: New test.
> ---
>  gcc/ChangeLog.openacc                              |  5 ++
>  gcc/c/ChangeLog.openacc                            |  5 ++
>  gcc/c/c-parser.c                                   |  4 +-
>  gcc/cp/ChangeLog.openacc                           |  5 ++
>  gcc/cp/parser.c                                    |  4 +-
>  gcc/fortran/ChangeLog.openacc                      |  5 ++
>  gcc/fortran/openmp.c                               |  4 +-
>  gcc/omp-expand.c                                   | 12 ++++-
>  gcc/testsuite/ChangeLog.openacc                    |  6 +++
>  gcc/testsuite/c-c++-common/goacc/host_data-1.c     | 28 +++++++++++-
>  gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 12 ++++-
>  include/ChangeLog.openacc                          |  4 ++
>  include/gomp-constants.h                           |  2 +
>  libgomp/ChangeLog.openacc                          | 10 ++++
>  libgomp/libgomp.h                                  |  3 ++
>  libgomp/oacc-parallel.c                            | 11 +++--
>  libgomp/target.c                                   |  3 ++
>  .../libgomp.oacc-c-c++-common/host_data-6.c        | 53 ++++++++++++++++++++++
>  18 files changed, 167 insertions(+), 9 deletions(-)
>  create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
> 
> diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
> index 6a51b1e..66eba7b 100644
> --- a/gcc/ChangeLog.openacc
> +++ b/gcc/ChangeLog.openacc
> @@ -1,3 +1,8 @@
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* omp-expand.c (expand_omp_target): Handle if_present flag on
> +	OpenACC host_data construct.
> +
>  2018-12-21  Thomas Schwinge  <thomas@codesourcery.com>
>  
>  	* omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP
> diff --git a/gcc/c/ChangeLog.openacc b/gcc/c/ChangeLog.openacc
> index 10c00e5..e607ea8 100644
> --- a/gcc/c/ChangeLog.openacc
> +++ b/gcc/c/ChangeLog.openacc
> @@ -1,3 +1,8 @@
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
> +	and PRAGMA_OACC_CLAUSE_IF_PRESENT.
> +
>  2018-12-20  Julian Brown  <julian@codesourcery.com>
>  	    Maciej W. Rozycki  <macro@codesourcery.com>
>  
> diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
> index a352d54..2bc4f45 100644
> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -14876,7 +14876,9 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
>  */
>  
>  #define OACC_HOST_DATA_CLAUSE_MASK					\
> -	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
> +        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE)          \
> +        | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)                  \
> +        | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) )
>  
>  static tree
>  c_parser_oacc_host_data (location_t loc, c_parser *parser, bool *if_p)
> diff --git a/gcc/cp/ChangeLog.openacc b/gcc/cp/ChangeLog.openacc
> index 37b0028..76889c7 100644
> --- a/gcc/cp/ChangeLog.openacc
> +++ b/gcc/cp/ChangeLog.openacc
> @@ -1,3 +1,8 @@
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
> +	and PRAGMA_OACC_CLAUSE_IF_PRESENT.
> +
>  2018-12-20  Julian Brown  <julian@codesourcery.com>
>  	    Maciej W. Rozycki  <macro@codesourcery.com>
>  
> diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
> index 083700b..38b0a6d 100644
> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
> @@ -36973,7 +36973,9 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
>    structured-block  */
>  
>  #define OACC_HOST_DATA_CLAUSE_MASK					\
> -  ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
> +  ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE)                \
> +  | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)                        \
> +  | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) )
>  
>  static tree
>  cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
> diff --git a/gcc/fortran/ChangeLog.openacc b/gcc/fortran/ChangeLog.openacc
> index a369219..306871e 100644
> --- a/gcc/fortran/ChangeLog.openacc
> +++ b/gcc/fortran/ChangeLog.openacc
> @@ -1,3 +1,8 @@
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
> +	OMP_CLAUSE_IF_PRESENT.
> +
>  2018-12-20  Julian Brown  <julian@codesourcery.com>
>  	    Maciej W. Rozycki  <macro@codesourcery.com>
>  
> diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
> index 350f4b1..4273dee 100644
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -2107,7 +2107,9 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
>     | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT	\
>     | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_ATTACH)
>  #define OACC_HOST_DATA_CLAUSES \
> -  (omp_mask (OMP_CLAUSE_USE_DEVICE))
> +  (omp_mask (OMP_CLAUSE_USE_DEVICE)                                     \
> +   | OMP_CLAUSE_IF                                                      \
> +   | OMP_CLAUSE_IF_PRESENT)
>  #define OACC_LOOP_CLAUSES \
>    (omp_mask (OMP_CLAUSE_COLLAPSE)					\
>     | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR		\
> diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
> index ea264da..42c4910 100644
> --- a/gcc/omp-expand.c
> +++ b/gcc/omp-expand.c
> @@ -7254,7 +7254,17 @@ expand_omp_target (struct omp_region *region)
>    if (is_gimple_omp_oacc (entry_stmt))
>      {
>        /* By default, no GOACC_FLAGs are set.  */
> -      goacc_flags = integer_zero_node;
> +      int goacc_flags_i = 0;
> +
> +      if (start_ix != BUILT_IN_GOACC_UPDATE
> +	  && omp_find_clause (clauses, OMP_CLAUSE_IF_PRESENT))
> +	{
> +	  gcc_checking_assert (gimple_omp_target_kind (entry_stmt)
> +			       == GF_OMP_TARGET_KIND_OACC_HOST_DATA);
> +	  goacc_flags_i |= GOACC_FLAG_HOST_DATA_IF_PRESENT;
> +	}
> +
> +      goacc_flags = build_int_cst (integer_type_node, goacc_flags_i);
>      }
>    else
>      {
> diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
> index 473eb9d..2e4bd3d 100644
> --- a/gcc/testsuite/ChangeLog.openacc
> +++ b/gcc/testsuite/ChangeLog.openacc
> @@ -1,5 +1,11 @@
>  2018-12-21  Gergö Barany  <gergo@codesourcery.com>
>  
> +	* c-c++-common/goacc/host_data-1.c: Add tests of if and if_present
> +	clauses on host_data.
> +	* gfortran.dg/goacc/host_data-tree.f95: Likewise.
> +
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
>  	* c-c++-common/goacc/nested-reductions-fail.c: Renamed to...
>  	* c-c++-common/goacc/nested-reductions-parallel-fail.c: ...this file,
>  	with kernels tests...
> diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-1.c b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
> index 0c7a857..658b7a6 100644
> --- a/gcc/testsuite/c-c++-common/goacc/host_data-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
> @@ -7,6 +7,9 @@ f (void)
>  {
>  #pragma acc host_data use_device(v1)
>    ;
> +
> +#pragma acc host_data use_device(v1) if_present
> +  ;
>  }
>  
>  
> @@ -16,9 +19,32 @@ void
>  foo (float *x, float *y)
>  {
>    int n = 1 << 10;
> -#pragma acc data create(x[0:n]) copyout(y[0:n])
> +#pragma acc data create(x[0:n])
>    {
> +    bar (x, y);
> +
> +    /* This should fail at run time because y is not mapped.  */
>  #pragma acc host_data use_device(x,y)
>      bar (x, y);
> +
> +    /* y is still not mapped, but this should not fail at run time but
> +       continue execution with y remaining as the host address.  */
> +#pragma acc host_data use_device(x,y) if_present
> +    bar (x, y);
> +
> +#pragma acc data copyout(y[0:n])
> +    {
> +#pragma acc host_data use_device(x,y)
> +      bar (x, y);
> +
> +#pragma acc host_data use_device(x,y) if_present
> +      bar (x, y);
> +
> +#pragma acc host_data use_device(x,y) if(x != y)
> +      bar (x, y);
> +
> +#pragma acc host_data use_device(x,y) if_present if(x != y)
> +      bar (x, y);
> +    }
>    }
>  }
> diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
> index d44ca58..2ac1c0d 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
> +++ b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
> @@ -7,5 +7,15 @@ program test
>  
>    !$acc host_data use_device(p)
>    !$acc end host_data
> +
> +  !$acc host_data use_device(p) if (p == 42)
> +  !$acc end host_data
> +
> +  !$acc host_data use_device(p) if_present if (p == 43)
> +  !$acc end host_data
>  end program test
> -! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 3 "original" } }
> +! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 42;" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\)" 2 "original" } }
> +! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 43;" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\) if_present" 1 "original" } }
> diff --git a/include/ChangeLog.openacc b/include/ChangeLog.openacc
> index aa583ea..82058e7 100644
> --- a/include/ChangeLog.openacc
> +++ b/include/ChangeLog.openacc
> @@ -1,3 +1,7 @@
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant.
> +
>  2018-12-21  Thomas Schwinge  <thomas@codesourcery.com>
>  
>  	* gomp-constants.h (GOACC_FLAG_HOST_FALLBACK)
> diff --git a/include/gomp-constants.h b/include/gomp-constants.h
> index b5d8441..953df8f 100644
> --- a/include/gomp-constants.h
> +++ b/include/gomp-constants.h
> @@ -257,6 +257,8 @@ enum gomp_map_kind
>  
>  /* Force host fallback execution.  */
>  #define GOACC_FLAG_HOST_FALLBACK	(1 << 0)
> +/* "if_present" semantics for OpenACC "host_data" constructs.  */
> +#define GOACC_FLAG_HOST_DATA_IF_PRESENT	(1 << 1)
>  
>  /* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted
>     bitmask.  */
> diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc
> index 04cea5f..7b9e2c5 100644
> --- a/libgomp/ChangeLog.openacc
> +++ b/libgomp/ChangeLog.openacc
> @@ -1,3 +1,13 @@
> +2018-12-21  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* libgomp.h (enum gomp_map_vars_kind): Add
> +	GOMP_MAP_VARS_OPENACC_IF_PRESENT.
> +	* oacc-parallel.c (GOACC_data_start): Handle
> +	GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
> +	* target.c (gomp_map_vars_async): Handle
> +	GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.
> +	* testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.
> +
>  2018-12-21  Thomas Schwinge  <thomas@codesourcery.com>
>  
>  	* oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel)
> diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
> index 64895c5..11948d5 100644
> --- a/libgomp/libgomp.h
> +++ b/libgomp/libgomp.h
> @@ -1024,6 +1024,9 @@ struct gomp_device_descr
>  enum gomp_map_vars_kind
>  {
>    GOMP_MAP_VARS_OPENACC,
> +  /* Like "GOMP_MAP_VARS_OPENACC", but with "GOACC_FLAG_HOST_DATA_IF_PRESENT"
> +     semantics.  */
> +  GOMP_MAP_VARS_OPENACC_IF_PRESENT,
>    GOMP_MAP_VARS_OPENACC_ENTER_DATA,
>    GOMP_MAP_VARS_TARGET,
>    GOMP_MAP_VARS_DATA,
> diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
> index 0b5f41a..3da87a1 100644
> --- a/libgomp/oacc-parallel.c
> +++ b/libgomp/oacc-parallel.c
> @@ -585,6 +585,12 @@ GOACC_data_start (int flags_m, size_t mapnum,
>  
>    handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
>  
> +  enum gomp_map_vars_kind pragma_kind;
> +  if (flags & GOACC_FLAG_HOST_DATA_IF_PRESENT)
> +    pragma_kind = GOMP_MAP_VARS_OPENACC_IF_PRESENT;
> +  else
> +    pragma_kind = GOMP_MAP_VARS_OPENACC;
> +
>    /* Host fallback or 'do nothing'.  */
>    if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
>        || (flags & GOACC_FLAG_HOST_FALLBACK))
> @@ -592,8 +598,7 @@ GOACC_data_start (int flags_m, size_t mapnum,
>        //TODO
>        prof_info.device_type = acc_device_host;
>        api_info.device_type = prof_info.device_type;
> -      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
> -			   GOMP_MAP_VARS_OPENACC);
> +      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, pragma_kind);
>        tgt->prev = thr->mapped_data;
>        thr->mapped_data = tgt;
>        goto out;
> @@ -601,7 +606,7 @@ GOACC_data_start (int flags_m, size_t mapnum,
>  
>    gomp_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
>    tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
> -		       GOMP_MAP_VARS_OPENACC);
> +		       pragma_kind);
>    gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
>    tgt->prev = thr->mapped_data;
>    thr->mapped_data = tgt;
> diff --git a/libgomp/target.c b/libgomp/target.c
> index 0594405..bdfd640 100644
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -1137,6 +1137,9 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
>  	  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
>  	  if (n == NULL)
>  	    {
> +              if (pragma_kind == GOMP_MAP_VARS_OPENACC_IF_PRESENT)
> +                /* No error, continue using the host address.  */
> +                continue;
>  	      gomp_mutex_unlock (&devicep->lock);
>  	      gomp_fatal ("use_device_ptr pointer wasn't mapped");
>  	    }
> 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..c5744fe
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
> @@ -0,0 +1,53 @@
> +/* Test if, if_present clauses on host_data construct.  */
> +
> +#include <assert.h>
> +#include <stdint.h>
> +
> +void
> +foo (float *p, intptr_t host_p, int shared_mem_p, int cond)
> +{
> +  assert (p == (float *) host_p);
> +
> +#pragma acc data copyin(host_p)
> +  {
> +#pragma acc host_data use_device(p) if_present
> +    /* p not mapped yet, so it will be equal to the host pointer.  */
> +    assert (p == (float *) host_p);
> +
> +#pragma acc data copy(p[0:100])
> +    {
> +      /* Not inside a host_data construct, so p is still the host pointer.  */
> +      assert (p == (float *) host_p);
> +
> +      if (!shared_mem_p)
> +      {
> +#pragma acc host_data use_device(p)
> +        /* The device address is different from the host address.  */
> +        assert (p != (float *) host_p);
> +
> +#pragma acc host_data use_device(p) if_present
> +        /* p is present now, so this is the same as above.  */
> +        assert (p != (float *) host_p);
> +      }
> +
> +#pragma acc host_data use_device(p) if(cond)
> +      /* p is the device pointer iff cond is true and device memory is
> +         separate from host memory.  */
> +      assert ((p != (float *) host_p) == (cond && !shared_mem_p));
> +    }
> +  }
> +}
> +
> +int
> +main (void)
> +{
> +  float arr[100];
> +  int shared_mem_p = 0;
> +#if ACC_MEM_SHARED
> +  shared_mem_p = 1;
> +#endif
> +  foo (arr, (intptr_t) arr, shared_mem_p, 0);
> +  foo (arr, (intptr_t) arr, shared_mem_p, 1);
> +
> +  return 0;
> +}
> -- 
> 2.8.1
>