From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from us-smtp-delivery-124.mimecast.com (us-smtp-delivery-124.mimecast.com [170.10.129.124]) by sourceware.org (Postfix) with ESMTPS id 13DBC383820F for ; Mon, 13 Jun 2022 12:06:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 13DBC383820F Received: from mimecast-mx02.redhat.com (mimecast-mx02.redhat.com [66.187.233.88]) by relay.mimecast.com with ESMTP with STARTTLS (version=TLSv1.2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id us-mta-505-wxZE-RvcPzirDZTAektgbQ-1; Mon, 13 Jun 2022 08:06:44 -0400 X-MC-Unique: wxZE-RvcPzirDZTAektgbQ-1 Received: from smtp.corp.redhat.com (int-mx03.intmail.prod.int.rdu2.redhat.com [10.11.54.3]) (using TLSv1.2 with cipher AECDH-AES256-SHA (256/256 bits)) (No client certificate requested) by mimecast-mx02.redhat.com (Postfix) with ESMTPS id 2E8B6100BABB; Mon, 13 Jun 2022 12:06:44 +0000 (UTC) Received: from tucnak.zalov.cz (unknown [10.39.192.11]) by smtp.corp.redhat.com (Postfix) with ESMTPS id A86501121319; Mon, 13 Jun 2022 12:06:43 +0000 (UTC) Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.17.1/8.17.1) with ESMTPS id 25DC6e2Z3471431 (version=TLSv1.3 cipher=TLS_AES_256_GCM_SHA384 bits=256 verify=NOT); Mon, 13 Jun 2022 14:06:41 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.17.1/8.17.1/Submit) id 25DC6ebw3471430; Mon, 13 Jun 2022 14:06:40 +0200 Date: Mon, 13 Jun 2022 14:06:39 +0200 From: Jakub Jelinek To: gcc-patches@gcc.gnu.org Cc: Tobias Burnus Subject: [committed] openmp: Conforming device numbers and omp_{initial,invalid}_device Message-ID: Reply-To: Jakub Jelinek MIME-Version: 1.0 X-Scanned-By: MIMEDefang 2.78 on 10.11.54.3 X-Mimecast-Spam-Score: 0 X-Mimecast-Originator: redhat.com Content-Type: text/plain; charset=us-ascii Content-Disposition: inline X-Spam-Status: No, score=-4.2 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_NONE, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 13 Jun 2022 12:06:50 -0000 Hi! 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.). Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk. 2022-06-13 Jakub Jelinek 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. --- gcc/omp-expand.cc.jj 2022-05-30 14:07:02.075305621 +0200 +++ gcc/omp-expand.cc 2022-06-10 15:49:44.528206376 +0200 @@ -9983,6 +9983,8 @@ expand_omp_target (struct omp_region *re 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 *re 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, "% not yet supported"); @@ -10021,7 +10036,8 @@ expand_omp_target (struct omp_region *re 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 *re 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 *re 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 *re 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); --- include/gomp-constants.h.jj 2022-05-25 11:06:59.585503316 +0200 +++ include/gomp-constants.h 2022-06-10 13:51:15.147519826 +0200 @@ -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) --- libgomp/omp.h.in.jj 2022-05-20 11:45:17.963742623 +0200 +++ libgomp/omp.h.in 2022-06-10 13:43:35.154224513 +0200 @@ -184,6 +184,12 @@ typedef enum omp_event_handle_t __GOMP_U __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 () --- libgomp/omp_lib.f90.in.jj 2022-05-20 11:45:17.963742623 +0200 +++ libgomp/omp_lib.f90.in 2022-06-10 13:43:43.561138515 +0200 @@ -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 --- libgomp/omp_lib.h.in.jj 2022-05-20 11:45:17.976742449 +0200 +++ libgomp/omp_lib.h.in 2022-06-10 13:43:51.790054368 +0200 @@ -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 --- libgomp/target.c.jj 2022-05-23 10:59:06.280590872 +0200 +++ libgomp/target.c 2022-06-10 16:43:43.045150123 +0200 @@ -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) (voi 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 **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 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 *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, 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 devic 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 d 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_ 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_ *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 ()) - 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; @@ -3830,13 +3829,7 @@ omp_target_associate_ptr (const void *ho 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 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 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; } --- libgomp/icv-device.c.jj 2022-01-11 23:11:23.889269089 +0100 +++ libgomp/icv-device.c 2022-06-10 17:17:32.144759265 +0200 @@ -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) --- libgomp/libgomp.texi.jj 2022-06-04 10:34:26.410504167 +0200 +++ libgomp/libgomp.texi 2022-06-13 13:36:08.031440481 +0200 @@ -403,7 +403,7 @@ The OpenMP 4.5 specification is fully su @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 --- libgomp/testsuite/libgomp.c/target-41.c.jj 2022-05-27 12:48:40.762483765 +0200 +++ libgomp/testsuite/libgomp.c/target-41.c 2022-06-10 17:01:37.923277206 +0200 @@ -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; } --- libgomp/testsuite/libgomp.c/target-45.c.jj 2022-06-10 17:05:09.901162819 +0200 +++ libgomp/testsuite/libgomp.c/target-45.c 2022-06-10 17:13:56.395911268 +0200 @@ -0,0 +1,19 @@ +/* { dg-shouldfail "omp_invalid_device" } */ + +#include + +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" } */ --- libgomp/testsuite/libgomp.c/target-46.c.jj 2022-06-10 17:12:06.552006904 +0200 +++ libgomp/testsuite/libgomp.c/target-46.c 2022-06-10 17:13:48.611988903 +0200 @@ -0,0 +1,20 @@ +/* { dg-shouldfail "omp_invalid_device" } */ + +#include + +void +foo (void) +{ +} + +volatile int dev = omp_invalid_device; + +int +main () +{ + #pragma omp target device (dev) + foo (); + return 0; +} + +/* { dg-output "omp_invalid_device" } */ --- libgomp/testsuite/libgomp.c/target-47.c.jj 2022-06-10 17:13:02.465449198 +0200 +++ libgomp/testsuite/libgomp.c/target-47.c 2022-06-10 17:13:41.471060132 +0200 @@ -0,0 +1,19 @@ +/* { dg-shouldfail "omp_invalid_device" } */ + +#include + +void +foo (void) +{ +} + +int +main () +{ + omp_set_default_device (omp_invalid_device); + #pragma omp target + foo (); + return 0; +} + +/* { dg-output "omp_invalid_device" } */ --- libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c.jj 2022-05-23 21:44:48.950848384 +0200 +++ libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c 2022-06-13 13:10:56.471535852 +0200 @@ -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)) --- libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90.jj 2022-05-23 21:44:48.954848343 +0200 +++ libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 2022-06-13 13:12:08.133819977 +0200 @@ -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 Jakub