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.
>From 6d719cc2bcfa8f7ed8cb59e753e44aab6bf634fb Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@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 <tho...@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 <ge...@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 <tho...@codesourcery.com>
+
+ * gomp-constants.h (GOACC_FLAG_HOST_FALLBACK)
+ (GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define.
+
2018-12-20 Julian Brown <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@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 <tho...@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 <ge...@codesourcery.com>
Thomas Schwinge <tho...@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?= <ge...@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 <ge...@codesourcery.com>
+
+ * omp-expand.c (expand_omp_target): Handle if_present flag on
+ OpenACC host_data construct.
+
2018-12-21 Thomas Schwinge <tho...@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 <ge...@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 <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@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 <ge...@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 <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@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 <ge...@codesourcery.com>
+
+ * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
+ OMP_CLAUSE_IF_PRESENT.
+
2018-12-20 Julian Brown <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@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 <ge...@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 <ge...@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 <ge...@codesourcery.com>
+
+ * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant.
+
2018-12-21 Thomas Schwinge <tho...@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 <ge...@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 <tho...@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
