mirror of
git://gcc.gnu.org/git/gcc.git
synced 2024-11-28 20:20:51 +08:00
openmp: Conforming device numbers and omp_{initial,invalid}_device
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.
This commit is contained in:
parent
3b598848f6
commit
1158fe4340
@ -9983,6 +9983,8 @@ expand_omp_target (struct omp_region *region)
|
||||
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. */
|
||||
@ -9994,6 +9996,19 @@ expand_omp_target (struct omp_region *region)
|
||||
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");
|
||||
@ -10021,7 +10036,8 @@ expand_omp_target (struct omp_region *region)
|
||||
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)
|
||||
{
|
||||
@ -10029,20 +10045,13 @@ expand_omp_target (struct omp_region *region)
|
||||
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
|
||||
@ -10067,6 +10076,7 @@ expand_omp_target (struct omp_region *region)
|
||||
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))
|
||||
@ -10099,6 +10109,50 @@ expand_omp_target (struct omp_region *region)
|
||||
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);
|
||||
|
@ -233,8 +233,19 @@ enum gomp_map_kind
|
||||
#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)
|
||||
|
@ -32,7 +32,7 @@ void
|
||||
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)
|
||||
|
@ -403,7 +403,7 @@ The OpenMP 4.5 specification is fully supported.
|
||||
@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
|
||||
|
@ -184,6 +184,12 @@ typedef enum omp_event_handle_t __GOMP_UINTPTR_T_ENUM
|
||||
__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 ()
|
||||
|
@ -168,6 +168,8 @@
|
||||
parameter :: omp_high_bw_mem_space = 3
|
||||
integer (omp_memspace_handle_kind), &
|
||||
parameter :: omp_low_lat_mem_space = 4
|
||||
integer, parameter :: omp_initial_device = -1
|
||||
integer, parameter :: omp_invalid_device = -4
|
||||
|
||||
type omp_alloctrait
|
||||
integer (kind=omp_alloctrait_key_kind) key
|
||||
|
@ -174,6 +174,9 @@
|
||||
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
|
||||
|
127
libgomp/target.c
127
libgomp/target.c
@ -126,18 +126,31 @@ gomp_get_num_devices (void)
|
||||
}
|
||||
|
||||
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");
|
||||
@ -2588,7 +2601,7 @@ GOMP_target (int device, void (*fn) (void *), 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);
|
||||
|
||||
void *fn_addr;
|
||||
if (devicep == NULL
|
||||
@ -2647,7 +2660,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
|
||||
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;
|
||||
|
||||
@ -2805,7 +2818,7 @@ void
|
||||
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)
|
||||
@ -2824,7 +2837,7 @@ void
|
||||
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)
|
||||
@ -2855,7 +2868,7 @@ void
|
||||
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)
|
||||
@ -2870,7 +2883,7 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
|
||||
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
|
||||
@ -3063,7 +3076,7 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
|
||||
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
|
||||
@ -3296,13 +3309,11 @@ GOMP_teams4 (unsigned int num_teams_low, unsigned int num_teams_high,
|
||||
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;
|
||||
|
||||
@ -3319,20 +3330,15 @@ omp_target_alloc (size_t size, int device_num)
|
||||
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)
|
||||
@ -3350,19 +3356,17 @@ omp_target_free (void *device_ptr, int device_num)
|
||||
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;
|
||||
@ -3384,12 +3388,11 @@ omp_target_memcpy_check (int dst_device_num, int src_device_num,
|
||||
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;
|
||||
|
||||
@ -3398,12 +3401,10 @@ omp_target_memcpy_check (int dst_device_num, int src_device_num,
|
||||
*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;
|
||||
|
||||
@ -3767,13 +3768,11 @@ int
|
||||
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 ())
|
||||
if (device_num == omp_initial_device
|
||||
|| 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;
|
||||
|
||||
@ -3830,13 +3829,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
|
||||
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;
|
||||
|
||||
@ -3872,13 +3865,11 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
|
||||
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;
|
||||
|
||||
@ -3910,13 +3901,11 @@ omp_get_mapped_ptr (const void *ptr, int device_num)
|
||||
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;
|
||||
|
||||
@ -3929,10 +3918,14 @@ int
|
||||
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;
|
||||
}
|
||||
|
@ -17,7 +17,10 @@ main ()
|
||||
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))
|
||||
|
@ -18,16 +18,18 @@ main ()
|
||||
{
|
||||
/* 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;
|
||||
}
|
||||
|
19
libgomp/testsuite/libgomp.c/target-45.c
Normal file
19
libgomp/testsuite/libgomp.c/target-45.c
Normal file
@ -0,0 +1,19 @@
|
||||
/* { 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" } */
|
20
libgomp/testsuite/libgomp.c/target-46.c
Normal file
20
libgomp/testsuite/libgomp.c/target-46.c
Normal file
@ -0,0 +1,20 @@
|
||||
/* { 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" } */
|
19
libgomp/testsuite/libgomp.c/target-47.c
Normal file
19
libgomp/testsuite/libgomp.c/target-47.c
Normal file
@ -0,0 +1,19 @@
|
||||
/* { 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" } */
|
@ -19,12 +19,15 @@ program main
|
||||
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.
|
||||
@ -35,14 +38,14 @@ program main
|
||||
!$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
|
||||
|
Loading…
Reference in New Issue
Block a user