OpenMP 5.2 changed once more what device numbers are allowed.
In 5.1, valid device numbers were [0, omp_get_num_devices()].
5.2 makes also -1 valid (calls it omp_initial_device), which is equivalent
in behavior to omp_get_num_devices() number but has the advantage that it
is a constant. And it also introduces omp_invalid_device which is
also a constant with implementation defined value < -1. That value should
act like sNaN, any time any device construct (GOMP_target*) or OpenMP runtime
API routine is asked for such a device, the program is terminated.
And if OMP_TARGET_OFFLOAD=mandatory, all non-conforming device numbers (which
is all but [-1, omp_get_num_devices()] other than omp_invalid_device)
must be treated like omp_invalid_device.
For device constructs, we have a compatibility problem, we've historically
used 2 magic negative values to mean something special.
GOMP_DEVICE_ICV (-1) means device clause wasn't present, pick the
omp_get_default_device () number
GOMP_DEVICE_FALLBACK (-2) means the host device (this is used e.g. for
#pragma omp target if (cond)
where if cond is false, we pass -2
But 5.2 requires that omp_initial_device is -1 (there were discussions
about it, advantage of -1 is that one can say iterate over the
[-1, omp_get_num_devices()-1] range to get all devices starting with
the host/initial one.
And also, if user passes -2, unless it is omp_invalid_device, we need to
treat it like non-conforming with OMP_TARGET_OFFLOAD=mandatory.
So, the patch does on the compiler side some number remapping,
user_device_num >= -2U ? user_device_num - 1 : user_device_num.
This remapping is done at compile time if device clause has constant
argument, otherwise at runtime, and means that for user -1 (omp_initial_device)
we pass -2 to GOMP_* in the runtime library where it treats it like host
fallback, while -2 is remapped to -3 (one of the non-conforming device numbers,
for those it doesn't matter which one is which).
omp_invalid_device is then -4.
For the OpenMP device runtime APIs, no remapping is done.
This patch doesn't deal with the initial default-device-var for
OMP_TARGET_OFFLOAD=mandatory , the spec says that the inital ICV value
for that should in that case depend on whether there are any offloading
devices or not (if not, should be omp_invalid_device), but that means
we can't determine the number of devices lazily (and let libraries have the
possibility to register their offloading data etc.).
2022-06-13 Jakub Jelinek <jakub@redhat.com>
gcc/
* omp-expand.cc (expand_omp_target): Remap user provided
device clause arguments, -1 to -2 and -2 to -3, either
at compile time if constant, or at runtime.
include/
* gomp-constants.h (GOMP_DEVICE_INVALID): Define.
libgomp/
* omp.h.in (omp_initial_device, omp_invalid_device): New enumerators.
* omp_lib.f90.in (omp_initial_device, omp_invalid_device): New
parameters.
* omp_lib.h.in (omp_initial_device, omp_invalid_device): Likewise.
* target.c (resolve_device): Add remapped argument, handle
GOMP_DEVICE_ICV only if remapped is true (and clear remapped),
for negative values, treat GOMP_DEVICE_FALLBACK as fallback only
if remapped, otherwise treat omp_initial_device that way. For
omp_invalid_device, always emit gomp_fatal, even when
OMP_TARGET_OFFLOAD isn't mandatory.
(GOMP_target, GOMP_target_ext, GOMP_target_data, GOMP_target_data_ext,
GOMP_target_update, GOMP_target_update_ext,
GOMP_target_enter_exit_data): Pass true as remapped argument to
resolve_device.
(omp_target_alloc, omp_target_free, omp_target_is_present,
omp_target_memcpy_check, omp_target_associate_ptr,
omp_target_disassociate_ptr, omp_get_mapped_ptr,
omp_target_is_accessible): Pass false as remapped argument to
resolve_device. Treat omp_initial_device the same as
gomp_get_num_devices (). Don't bypass resolve_device calls if
device_num is negative.
(omp_pause_resource): Treat omp_initial_device the same as
gomp_get_num_devices (). Call resolve_device.
* icv-device.c (omp_set_default_device): Always set to device_num
even when it is negative.
* libgomp.texi: Document that Conforming device numbers,
omp_initial_device and omp_invalid_device is implemented.
* testsuite/libgomp.c/target-41.c (main): Add test with
omp_initial_device.
* testsuite/libgomp.c/target-45.c: New test.
* testsuite/libgomp.c/target-46.c: New test.
* testsuite/libgomp.c/target-47.c: New test.
* testsuite/libgomp.c-c++-common/target-is-accessible-1.c (main): Add
test with omp_initial_device. Use -5 instead of -1 for negative value
test.
* testsuite/libgomp.fortran/target-is-accessible-1.f90 (main):
Likewise. Reorder stop numbers.
(cherry picked from commit
1158fe43407568f20415b16575ddbfff216bf8b6)
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-13 Jakub Jelinek <jakub@redhat.com>
+
+ * omp-expand.cc (expand_omp_target): Remap user provided
+ device clause arguments, -1 to -2 and -2 to -3, either
+ at compile time if constant, or at runtime.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
tree device = NULL_TREE;
location_t device_loc = UNKNOWN_LOCATION;
tree goacc_flags = NULL_TREE;
+ bool need_device_adjustment = false;
+ gimple_stmt_iterator adj_gsi;
if (is_gimple_omp_oacc (entry_stmt))
{
/* By default, no GOACC_FLAGs are set. */
if (c)
{
device = OMP_CLAUSE_DEVICE_ID (c);
+ /* Ensure 'device' is of the correct type. */
+ device = fold_convert_loc (device_loc, integer_type_node, device);
+ if (TREE_CODE (device) == INTEGER_CST)
+ {
+ if (wi::to_wide (device) == GOMP_DEVICE_ICV)
+ device = build_int_cst (integer_type_node,
+ GOMP_DEVICE_HOST_FALLBACK);
+ else if (wi::to_wide (device) == GOMP_DEVICE_HOST_FALLBACK)
+ device = build_int_cst (integer_type_node,
+ GOMP_DEVICE_HOST_FALLBACK - 1);
+ }
+ else
+ need_device_adjustment = true;
device_loc = OMP_CLAUSE_LOCATION (c);
if (OMP_CLAUSE_DEVICE_ANCESTOR (c))
sorry_at (device_loc, "%<ancestor%> not yet supported");
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)
+ OpenACC: goacc_flags = (cond ? goacc_flags
+ : goacc_flags | GOACC_FLAG_HOST_FALLBACK)
OpenMP: device = (cond ? device : GOMP_DEVICE_HOST_FALLBACK) */
if (cond)
{
if (is_gimple_omp_oacc (entry_stmt))
tp = &goacc_flags;
else
- {
- /* Ensure 'device' is of the correct type. */
- device = fold_convert_loc (device_loc, integer_type_node, device);
-
- tp = &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 (*tp));
+ tree tmp_var = create_tmp_var (TREE_TYPE (*tp));
if (offloaded)
e = split_block_after_labels (new_bb);
else
gsi = gsi_start_bb (then_bb);
stmt = gimple_build_assign (tmp_var, *tp);
gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ adj_gsi = gsi;
gsi = gsi_start_bb (else_bb);
if (is_gimple_omp_oacc (entry_stmt))
if (device != NULL_TREE)
device = force_gimple_operand_gsi (&gsi, device, true, NULL_TREE,
true, GSI_SAME_STMT);
+ if (need_device_adjustment)
+ {
+ tree tmp_var = create_tmp_var (TREE_TYPE (device));
+ stmt = gimple_build_assign (tmp_var, device);
+ gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+ adj_gsi = gsi_for_stmt (stmt);
+ device = tmp_var;
+ }
+ }
+
+ if (need_device_adjustment)
+ {
+ tree uns = fold_convert (unsigned_type_node, device);
+ uns = force_gimple_operand_gsi (&adj_gsi, uns, true, NULL_TREE,
+ false, GSI_CONTINUE_LINKING);
+ edge e = split_block (gsi_bb (adj_gsi), gsi_stmt (adj_gsi));
+ basic_block cond_bb = e->src;
+ basic_block else_bb = e->dest;
+ if (gsi_bb (adj_gsi) == new_bb)
+ {
+ new_bb = else_bb;
+ gsi = gsi_last_nondebug_bb (new_bb);
+ }
+
+ basic_block then_bb = create_empty_bb (cond_bb);
+ set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+
+ cond = build2 (GT_EXPR, boolean_type_node, uns,
+ build_int_cst (unsigned_type_node,
+ GOMP_DEVICE_HOST_FALLBACK - 1));
+ stmt = gimple_build_cond_empty (cond);
+ adj_gsi = gsi_last_bb (cond_bb);
+ gsi_insert_after (&adj_gsi, stmt, GSI_CONTINUE_LINKING);
+
+ adj_gsi = gsi_start_bb (then_bb);
+ tree add = build2 (PLUS_EXPR, integer_type_node, device,
+ build_int_cst (integer_type_node, -1));
+ stmt = gimple_build_assign (device, add);
+ gsi_insert_after (&adj_gsi, stmt, GSI_CONTINUE_LINKING);
+
+ make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+ e->flags = EDGE_FALSE_VALUE;
+ add_bb_to_loop (then_bb, cond_bb->loop_father);
+ make_edge (then_bb, else_bb, EDGE_FALLTHRU);
}
t = gimple_omp_target_data_arg (entry_stmt);
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-13 Jakub Jelinek <jakub@redhat.com>
+
+ * gomp-constants.h (GOMP_DEVICE_INVALID): Define.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
#define GOMP_DEVICE_HSA 7
#define GOMP_DEVICE_GCN 8
+/* We have a compatibility issue. OpenMP 5.2 introduced
+ omp_initial_device with value of -1 which clashes with our
+ GOMP_DEVICE_ICV, so we need to remap user supplied device
+ ids, -1 (aka omp_initial_device) to GOMP_DEVICE_HOST_FALLBACK,
+ and -2 (one of many non-conforming device numbers, but with
+ OMP_TARGET_OFFLOAD=mandatory needs to be treated a
+ omp_invalid_device) to -3 (so that for dev_num >= -2U we can
+ subtract 1). -4 is then what we use for omp_invalid_device,
+ which unlike the other non-conforming device numbers results
+ in fatal error regardless of OMP_TARGET_OFFLOAD. */
#define GOMP_DEVICE_ICV -1
#define GOMP_DEVICE_HOST_FALLBACK -2
+#define GOMP_DEVICE_INVALID -4
/* GOMP_task/GOMP_taskloop* flags argument. */
#define GOMP_TASK_FLAG_UNTIED (1 << 0)
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-13 Jakub Jelinek <jakub@redhat.com>
+
+ * omp.h.in (omp_initial_device, omp_invalid_device): New enumerators.
+ * omp_lib.f90.in (omp_initial_device, omp_invalid_device): New
+ parameters.
+ * omp_lib.h.in (omp_initial_device, omp_invalid_device): Likewise.
+ * target.c (resolve_device): Add remapped argument, handle
+ GOMP_DEVICE_ICV only if remapped is true (and clear remapped),
+ for negative values, treat GOMP_DEVICE_FALLBACK as fallback only
+ if remapped, otherwise treat omp_initial_device that way. For
+ omp_invalid_device, always emit gomp_fatal, even when
+ OMP_TARGET_OFFLOAD isn't mandatory.
+ (GOMP_target, GOMP_target_ext, GOMP_target_data, GOMP_target_data_ext,
+ GOMP_target_update, GOMP_target_update_ext,
+ GOMP_target_enter_exit_data): Pass true as remapped argument to
+ resolve_device.
+ (omp_target_alloc, omp_target_free, omp_target_is_present,
+ omp_target_memcpy_check, omp_target_associate_ptr,
+ omp_target_disassociate_ptr, omp_get_mapped_ptr,
+ omp_target_is_accessible): Pass false as remapped argument to
+ resolve_device. Treat omp_initial_device the same as
+ gomp_get_num_devices (). Don't bypass resolve_device calls if
+ device_num is negative.
+ (omp_pause_resource): Treat omp_initial_device the same as
+ gomp_get_num_devices (). Call resolve_device.
+ * icv-device.c (omp_set_default_device): Always set to device_num
+ even when it is negative.
+ * libgomp.texi: Document that Conforming device numbers,
+ omp_initial_device and omp_invalid_device is implemented.
+ * testsuite/libgomp.c/target-41.c (main): Add test with
+ omp_initial_device.
+ * testsuite/libgomp.c/target-45.c: New test.
+ * testsuite/libgomp.c/target-46.c: New test.
+ * testsuite/libgomp.c/target-47.c: New test.
+ * testsuite/libgomp.c-c++-common/target-is-accessible-1.c (main): Add
+ test with omp_initial_device. Use -5 instead of -1 for negative value
+ test.
+ * testsuite/libgomp.fortran/target-is-accessible-1.f90 (main):
+ Likewise. Reorder stop numbers.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
omp_set_default_device (int device_num)
{
struct gomp_task_icv *icv = gomp_icv (true);
- icv->default_device_var = device_num >= 0 ? device_num : 0;
+ icv->default_device_var = device_num;
}
ialias (omp_set_default_device)
@headitem Description @tab Status @tab Comments
@item For Fortran, optional comma between directive and clause @tab N @tab
@item Conforming device numbers and @code{omp_initial_device} and
- @code{omp_invalid_device} enum/PARAMETER @tab N @tab
+ @code{omp_invalid_device} enum/PARAMETER @tab Y @tab
@item Initial value of @emph{default-device-var} ICV with
@code{OMP_TARGET_OFFLOAD=mandatory} @tab N @tab
@item @emph{interop_types} in any position of the modifier list for the @code{init} clause
__omp_event_handle_t_max__ = __UINTPTR_MAX__
} omp_event_handle_t;
+enum
+{
+ omp_initial_device = -1,
+ omp_invalid_device = -4
+};
+
#ifdef __cplusplus
extern "C" {
# define __GOMP_NOTHROW throw ()
parameter :: omp_unified_shared_mem_space = 5
integer (omp_memspace_handle_kind), &
parameter :: omp_host_mem_space = 6
+ integer, parameter :: omp_initial_device = -1
+ integer, parameter :: omp_invalid_device = -4
type omp_alloctrait
integer (kind=omp_alloctrait_key_kind) key
parameter (omp_const_mem_space = 2)
parameter (omp_high_bw_mem_space = 3)
parameter (omp_low_lat_mem_space = 4)
+ integer omp_initial_device, omp_invalid_device
+ parameter (omp_initial_device = -1)
+ parameter (omp_invalid_device = -4)
type omp_alloctrait
integer (omp_alloctrait_key_kind) key
}
static struct gomp_device_descr *
-resolve_device (int device_id)
+resolve_device (int device_id, bool remapped)
{
- if (device_id == GOMP_DEVICE_ICV)
+ if (remapped && device_id == GOMP_DEVICE_ICV)
{
struct gomp_task_icv *icv = gomp_icv (false);
device_id = icv->default_device_var;
+ remapped = false;
}
- if (device_id < 0 || device_id >= gomp_get_num_devices ())
+ if (device_id < 0)
+ {
+ if (device_id == (remapped ? GOMP_DEVICE_HOST_FALLBACK
+ : omp_initial_device))
+ return NULL;
+ if (device_id == omp_invalid_device)
+ gomp_fatal ("omp_invalid_device encountered");
+ else if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY)
+ gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
+ "but device not found");
+
+ return NULL;
+ }
+ else if (device_id >= gomp_get_num_devices ())
{
if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
- && device_id != GOMP_DEVICE_HOST_FALLBACK
&& device_id != num_devices_openmp)
gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
"but device not found");
size_t mapnum, void **hostaddrs, size_t *sizes,
unsigned char *kinds)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
void *fn_addr;
if (devicep == NULL
void **hostaddrs, size_t *sizes, unsigned short *kinds,
unsigned int flags, void **depend, void **args)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
size_t tgt_align = 0, tgt_size = 0;
bool fpc_done = false;
GOMP_target_data (int device, const void *unused, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned char *kinds)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
size_t *sizes, unsigned short *kinds)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
GOMP_target_update (int device, const void *unused, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned char *kinds)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
size_t *sizes, unsigned short *kinds,
unsigned int flags, void **depend)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
/* If there are depend clauses, but nowait is not present,
block the parent task until the dependencies are resolved
size_t *sizes, unsigned short *kinds,
unsigned int flags, void **depend)
{
- struct gomp_device_descr *devicep = resolve_device (device);
+ struct gomp_device_descr *devicep = resolve_device (device, true);
/* If there are depend clauses, but nowait is not present,
block the parent task until the dependencies are resolved
void *
omp_target_alloc (size_t size, int device_num)
{
- if (device_num == gomp_get_num_devices ())
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
return malloc (size);
- if (device_num < 0)
- return NULL;
-
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
if (devicep == NULL)
return NULL;
void
omp_target_free (void *device_ptr, int device_num)
{
- if (device_ptr == NULL)
- return;
-
- if (device_num == gomp_get_num_devices ())
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
{
free (device_ptr);
return;
}
- if (device_num < 0)
- return;
-
- struct gomp_device_descr *devicep = resolve_device (device_num);
- if (devicep == NULL)
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
+ if (devicep == NULL || device_ptr == NULL)
return;
if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
if (device_num == gomp_get_num_devices ())
return malloc (size);
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, true);
if (devicep == NULL)
return NULL;
return;
}
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, true);
if (devicep == NULL)
return;
int
omp_target_is_present (const void *ptr, int device_num)
{
- if (ptr == NULL)
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
return 1;
- if (device_num == gomp_get_num_devices ())
- return 1;
-
- if (device_num < 0)
- return 0;
-
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
if (devicep == NULL)
return 0;
+ if (ptr == NULL)
+ return 1;
+
if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|| devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
return 1;
struct gomp_device_descr **dst_devicep,
struct gomp_device_descr **src_devicep)
{
- if (dst_device_num != gomp_get_num_devices ())
+ if (dst_device_num != gomp_get_num_devices ()
+ /* Above gomp_get_num_devices has to be called unconditionally. */
+ && dst_device_num != omp_initial_device)
{
- if (dst_device_num < 0)
- return EINVAL;
-
- *dst_devicep = resolve_device (dst_device_num);
+ *dst_devicep = resolve_device (dst_device_num, false);
if (*dst_devicep == NULL)
return EINVAL;
*dst_devicep = NULL;
}
- if (src_device_num != num_devices_openmp)
+ if (src_device_num != num_devices_openmp
+ && src_device_num != omp_initial_device)
{
- if (src_device_num < 0)
- return EINVAL;
-
- *src_devicep = resolve_device (src_device_num);
+ *src_devicep = resolve_device (src_device_num, false);
if (*src_devicep == NULL)
return EINVAL;
omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
size_t size, size_t device_offset, int device_num)
{
- if (device_num == gomp_get_num_devices ())
- return EINVAL;
-
- if (device_num < 0)
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
return EINVAL;
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
if (devicep == NULL)
return EINVAL;
int
omp_target_disassociate_ptr (const void *ptr, int device_num)
{
- if (device_num == gomp_get_num_devices ())
- return EINVAL;
-
- if (device_num < 0)
- return EINVAL;
-
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
if (devicep == NULL)
return EINVAL;
void *
omp_get_mapped_ptr (const void *ptr, int device_num)
{
- if (device_num < 0 || device_num > gomp_get_num_devices ())
- return NULL;
-
- if (device_num == omp_get_initial_device ())
+ if (device_num == omp_initial_device
+ || device_num == omp_get_initial_device ())
return (void *) ptr;
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
if (devicep == NULL)
return NULL;
int
omp_target_is_accessible (const void *ptr, size_t size, int device_num)
{
- if (device_num < 0 || device_num > gomp_get_num_devices ())
- return false;
-
- if (device_num == gomp_get_num_devices ())
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
return true;
- struct gomp_device_descr *devicep = resolve_device (device_num);
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
if (devicep == NULL)
return false;
omp_pause_resource (omp_pause_resource_t kind, int device_num)
{
(void) kind;
- if (device_num == gomp_get_num_devices ())
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
return gomp_pause_host ();
- if (device_num < 0 || device_num >= num_devices_openmp)
+
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
+ if (devicep == NULL)
return -1;
+
/* Do nothing for target devices for now. */
return 0;
}
result = false;
else
{
- struct gomp_device_descr *device = resolve_device (device_num);
+ struct gomp_device_descr *device = resolve_device (device_num, true);
if (device == NULL)
result = false;
else if (device->evaluate_device_func)
if (!omp_target_is_accessible (p, sizeof (int), id))
__builtin_abort ();
- if (omp_target_is_accessible (p, sizeof (int), -1))
+ if (!omp_target_is_accessible (p, sizeof (int), omp_initial_device))
+ __builtin_abort ();
+
+ if (omp_target_is_accessible (p, sizeof (int), -5))
__builtin_abort ();
if (omp_target_is_accessible (p, sizeof (int), n + 1))
{
/* OMP_TARGET_OFFLOAD=mandatory shouldn't fail for host fallback
if it is because the program explicitly asked for the host
- fallback through if(false) or omp_get_initial_device () as
- the device. */
+ fallback through if(false) or omp_get_initial_device () or
+ omp_initial_device as the device. */
#pragma omp target if (v)
foo ();
+ #pragma omp target device (omp_initial_device)
+ foo ();
#pragma omp target device (omp_get_initial_device ())
foo ();
omp_set_default_device (omp_get_initial_device ());
#pragma omp target
foo ();
- if (v != 3)
+ if (v != 4)
abort ();
return 0;
}
--- /dev/null
+/* { dg-shouldfail "omp_invalid_device" } */
+
+#include <omp.h>
+
+void
+foo (void)
+{
+}
+#pragma omp declare target enter (foo)
+
+int
+main ()
+{
+ #pragma omp target device (omp_invalid_device)
+ foo ();
+ return 0;
+}
+
+/* { dg-output "omp_invalid_device" } */
--- /dev/null
+/* { dg-shouldfail "omp_invalid_device" } */
+
+#include <omp.h>
+
+void
+foo (void)
+{
+}
+
+volatile int dev = omp_invalid_device;
+
+int
+main ()
+{
+ #pragma omp target device (dev)
+ foo ();
+ return 0;
+}
+
+/* { dg-output "omp_invalid_device" } */
--- /dev/null
+/* { dg-shouldfail "omp_invalid_device" } */
+
+#include <omp.h>
+
+void
+foo (void)
+{
+}
+
+int
+main ()
+{
+ omp_set_default_device (omp_invalid_device);
+ #pragma omp target
+ foo ();
+ return 0;
+}
+
+/* { dg-output "omp_invalid_device" } */
if (omp_target_is_accessible (p, c_sizeof (d), id) /= 1) &
stop 2
- if (omp_target_is_accessible (p, c_sizeof (d), -1) /= 0) &
+ if (omp_target_is_accessible (p, c_sizeof (d), omp_initial_device) /= 1) &
stop 3
- if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) &
+ if (omp_target_is_accessible (p, c_sizeof (d), -5) /= 0) &
stop 4
+ if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) &
+ stop 5
+
! Currently, a host pointer is accessible if the device supports shared
! memory or omp_target_is_accessible is executed on the host. This
! test case must be adapted when unified shared memory is avialable.
!$omp end target
if (omp_target_is_accessible (p, c_sizeof (d), d) /= shared_mem) &
- stop 5;
+ stop 6;
if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) &
- stop 6;
+ stop 7;
do i = 1, 128
if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) &
- stop 7;
+ stop 8;
end do
end do
projects / thirdparty / gcc.git / commitdiff
