From patchwork Tue Jun 13 18:44:39 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1794630 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4QgctS5jCvz20QH for ; Wed, 14 Jun 2023 04:45:04 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D34823858281 for ; Tue, 13 Jun 2023 18:45:01 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 106863858D38 for ; Tue, 13 Jun 2023 18:44:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 106863858D38 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.00,240,1681200000"; d="diff'?scan'208";a="8680220" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 13 Jun 2023 10:44:44 -0800 IronPort-SDR: QgzLl0rBi5UkzcVA9+BiU4caqqTFTRJMD2077nRWWc4jC4sbbsdsk/o2gr0QvM6O4jm7ku+nSu ZZ7SDsJji4gx7F8Q5S6VnlfsNc5UJ6YZpoTzFUXxbrfeRRxW2wCZI/68XNoLr43cJco1KqNQD0 CClVeEfLyCAMZtnN802dNnsmT9gFhmTSpjjK6kEdzxVuAcDSuIdSR5jQHeA0cLm00IxWyqY3hO lPPCqeo7EBziVEG3szPI5qZGIxwXXnGIzLZKpuksOsNGnDojgpVU7TUAy5u2rSBR+w2owLfJ1v V3c= Message-ID: <1487d7d4-8611-0d78-6bf2-9bffdd4daa64@codesourcery.com> Date: Tue, 13 Jun 2023 20:44:39 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.12.0 Content-Language: en-US To: gcc-patches From: Tobias Burnus Subject: [patch] OpenMP: Set default-device-var with OMP_TARGET_OFFLOAD=mandatory X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, 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: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" I intent to commit this tomorrow, unless there are comments. It does as it says (see commit log): It initializes default-device-var to the value using the algorithm described in OpenMP 5.2, which depends on whether OMP_TARGET_OFFLOAD=mandatory was set. NOTE: With -foffload=disable there is no binary code but still devices get found - such that default-device-var == 0 (= first nonhost device). Thus, in that case, libgomp runs the code on that device but as no binary data is available, host fallback is used. (Even if there would be executable code for another device on the system.) With mandatory, this unintended host fallback is detected and an error is diagnosed. One can argue whether keeping the devices makes sense (e.g. because in a dynamic library device code will be loaded later) or not (don't list if no code is available). Note that TR11 (future OpenMP 6.0) extends OMP_DEFAULT_DEVICE and adds OMP_AVAILABLE_DEVICES which permit a finer-grained control about the device, including OMP_DEFAULT_DEVICE=initial (and 'invalid') which the current scheme does not permit. (Well, there is OMP_TARGET_OFFLOAD=disabled, but that's a too big hammer.) Tobias PS: DejaGNU testing was done without offloading configured and with remote testing on a system having an offload device, which which does not support setting environment variables. Manual testing was done with offloading enabled and depending on the testcase, running on a system with and/or without offloading hardware. ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 OpenMP: Set default-device-var with OMP_TARGET_OFFLOAD=mandatory OMP_TARGET_OFFLOAD=mandatory handling was before inconsistent. Hence, in OpenMP 5.2 it was clarified/extended by having implications on the default-device-var; additionally, omp_initial_device and omp_invalid_device enum values/PARAMETERs were added; support for it was added in r13-1066-g1158fe43407568 including aborting for omp_invalid_device and non-conforming device numbers. Only the mandatory handling was missing. Namely, while the default-device-var is usually initialized to value 0, with 'mandatory' it must have the value 'omp_invalid_device' if and only if zero non-host devices are available. (The OMP_DEFAULT_DEVICE env var overrides this as it comes semantically after the initialization.) To achieve this, default-device-var is now initialized to MIN_INT. If there is no 'mandatory', it is set to 0 directly after env var parsing. Otherwise, it is updated in gomp_target_init to either 0 or omp_invalid_device. To ensure INT_MIN is never seen by the user, both the omp_get_default_device API routine and omp_display_env (user call and OMP_DISPLAY_ENV env var) call gomp_init_targets_once() in that case. libgomp/ChangeLog: * env.c (gomp_default_icv_values): Init default_device_var to an nonconforming value - INT_MIN. (initialize_env): After env-var parsing, set default_device_var to device 0 unless OMP_TARGET_OFFLOAD=mandatory. (omp_display_env): If default_device_var is INT_MIN, call gomp_init_targets_once. * icv-device.c (omp_get_default_device): Likewise. * libgomp.texi (OMP_DEFAULT_DEVICE): Update init description. (OpenMP 5.2 Impl. Status): Mark OMP_TARGET_OFFLOAD=mandatory as 'Y'. * target.c (resolve_device): Improve error message device-num < 0 with 'mandatory' and no no-host devices available. (gomp_target_init): Set default-device-var if INT_MIN. * testsuite/libgomp.c/target-48.c: New test. * testsuite/libgomp.c/target-49.c: New test. * testsuite/libgomp.c/target-50.c: New test. * testsuite/libgomp.c/target-51.c: New test. * testsuite/libgomp.c/target-52.c: New test. * testsuite/libgomp.c/target-53.c: New test. * testsuite/libgomp.c/target-54.c: New test. libgomp/env.c | 13 ++++++++-- libgomp/icv-device.c | 4 +++ libgomp/libgomp.texi | 4 ++- libgomp/target.c | 15 ++++++++++- libgomp/testsuite/libgomp.c/target-48.c | 31 +++++++++++++++++++++++ libgomp/testsuite/libgomp.c/target-49.c | 18 +++++++++++++ libgomp/testsuite/libgomp.c/target-50.c | 27 ++++++++++++++++++++ libgomp/testsuite/libgomp.c/target-50a.c | 43 ++++++++++++++++++++++++++++++++ libgomp/testsuite/libgomp.c/target-51.c | 24 ++++++++++++++++++ libgomp/testsuite/libgomp.c/target-52.c | 25 +++++++++++++++++++ libgomp/testsuite/libgomp.c/target-53.c | 22 ++++++++++++++++ libgomp/testsuite/libgomp.c/target-54.c | 20 +++++++++++++++ 12 files changed, 242 insertions(+), 4 deletions(-) diff --git a/libgomp/env.c b/libgomp/env.c index e7a035b593c..25c0211dda1 100644 --- a/libgomp/env.c +++ b/libgomp/env.c @@ -62,13 +62,14 @@ #include "secure_getenv.h" #include "environ.h" -/* Default values of ICVs according to the OpenMP standard. */ +/* Default values of ICVs according to the OpenMP standard, + except for default-device-var. */ const struct gomp_default_icv gomp_default_icv_values = { .nthreads_var = 1, .thread_limit_var = UINT_MAX, .run_sched_var = GFS_DYNAMIC, .run_sched_chunk_size = 1, - .default_device_var = 0, + .default_device_var = INT_MIN, .max_active_levels_var = 1, .bind_var = omp_proc_bind_false, .nteams_var = 0, @@ -1614,6 +1615,10 @@ omp_display_env (int verbose) struct gomp_icv_list *none = gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_NO_SUFFIX); + if (none->icvs.default_device_var == INT_MIN) + /* This implies OMP_TARGET_OFFLOAD=mandatory. */ + gomp_init_targets_once (); + fputs ("\nOPENMP DISPLAY ENVIRONMENT BEGIN\n", stderr); fputs (" _OPENMP = '201511'\n", stderr); @@ -2213,6 +2218,10 @@ initialize_env (void) gomp_global_icv.max_active_levels_var = gomp_supported_active_levels; } + if (gomp_global_icv.default_device_var == INT_MIN + && gomp_target_offload_var != GOMP_TARGET_OFFLOAD_MANDATORY) + none->icvs.default_device_var = gomp_global_icv.default_device_var = 0; + /* Process GOMP_* variables and dependencies between parsed ICVs. */ parse_int_secure ("GOMP_DEBUG", &gomp_debug_var, true); diff --git a/libgomp/icv-device.c b/libgomp/icv-device.c index a2bbedc672a..b48ea3b096c 100644 --- a/libgomp/icv-device.c +++ b/libgomp/icv-device.c @@ -27,6 +27,7 @@ expected to replace. */ #include "libgomp.h" +#include void omp_set_default_device (int device_num) @@ -41,6 +42,9 @@ int omp_get_default_device (void) { struct gomp_task_icv *icv = gomp_icv (false); + if (icv->default_device_var == INT_MIN) + /* This implies OMP_TARGET_OFFLOAD=mandatory. */ + gomp_init_targets_once (); return icv->default_device_var; } diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index a3d370a0fb3..21d3582a665 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -423,7 +423,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab @item Conforming device numbers and @code{omp_initial_device} and @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 + @code{OMP_TARGET_OFFLOAD=mandatory} @tab Y @tab @item @emph{interop_types} in any position of the modifier list for the @code{init} clause of the @code{interop} construct @tab N @tab @end multitable @@ -2006,6 +2006,8 @@ Set to choose the device which is used in a @code{target} region, unless the value is overridden by @code{omp_set_default_device} or by a @code{device} clause. The value shall be the nonnegative device number. If no device with the given device number exists, the code is executed on the host. If unset, +@env{OMP_TARGET_OFFLOAD} is @code{mandatory} and no non-host devices are +available, it is set to @code{omp_invalid_device}. Otherwise, if unset, device number 0 will be used. diff --git a/libgomp/target.c b/libgomp/target.c index e3c4121a09f..f1020fad601 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -150,7 +150,11 @@ resolve_device (int device_id, bool remapped) if (device_id == (remapped ? GOMP_DEVICE_HOST_FALLBACK : omp_initial_device)) return NULL; - if (device_id == omp_invalid_device) + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY + && gomp_get_num_devices () == 0) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY but only the host " + "device is available"); + else 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, " @@ -5184,6 +5188,15 @@ gomp_target_init (void) if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200) goacc_register (&devs[i]); } + if (gomp_global_icv.default_device_var == INT_MIN) + { + /* This implies OMP_TARGET_OFFLOAD=mandatory. */ + struct gomp_icv_list *none; + none = gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_NO_SUFFIX); + gomp_global_icv.default_device_var = (num_devs_openmp + ? 0 : omp_invalid_device); + none->icvs.default_device_var = gomp_global_icv.default_device_var; + } num_devices = num_devs; num_devices_openmp = num_devs_openmp; diff --git a/libgomp/testsuite/libgomp.c/target-48.c b/libgomp/testsuite/libgomp.c/target-48.c new file mode 100644 index 00000000000..8e95c1c3ac3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-48.c @@ -0,0 +1,31 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices; + omp_invalid_device == -4 with GCC. */ + +/* { dg-do run { target { ! offload_device } } } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '-4'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +#include + +int +main () +{ + if (omp_get_default_device () != omp_invalid_device) + __builtin_abort (); + + omp_set_default_device (omp_initial_device); + + /* The spec is a bit unclear whether the line above sets the device number + (a) to -1 (= omp_initial_device) or + (b) to omp_get_initial_device() == omp_get_num_devices(). Therefore, + we accept either value. */ + + if (omp_get_default_device() != omp_get_initial_device() + && omp_get_default_device() != omp_initial_device) + __builtin_abort (); + + omp_display_env (0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-49.c b/libgomp/testsuite/libgomp.c/target-49.c new file mode 100644 index 00000000000..970cb91d512 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-49.c @@ -0,0 +1,18 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices, + which is enforced by using -foffload=disable. */ + +/* { dg-do run } */ +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* See comment in target-50.c/target-50.c for why default-device-var can be '0'. */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '-4'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target { ! offload_device } } } */ +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target offload_device } } */ + +int +main () +{ + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-50.c b/libgomp/testsuite/libgomp.c/target-50.c new file mode 100644 index 00000000000..6f15569ee21 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-50.c @@ -0,0 +1,27 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices; + here with using -foffload=disable. + As default-device-var is set to 0 (= host in this case), it should not fail. */ + +/* Note that -foffload=disable will still find devices on the system and only + when trying to use them, it will fail as no binary data has been produced. + The "target offload_device" case is checked for in 'target-50a.c'. */ + +/* { dg-do run { target { ! offload_device } } } */ + +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "0" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-50a.c b/libgomp/testsuite/libgomp.c/target-50a.c new file mode 100644 index 00000000000..0835cb5bae3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-50a.c @@ -0,0 +1,43 @@ +/* Check OMP_TARGET_OFFLOAD on systems with non-host devices but no executable + code due to -foffload=disable. + + Note: While one might expect that -foffload=disable implies no non-host + devices, libgomp actually detects the devices and only fails when trying to + run as no executable code is availale for that device. + (Without MANDATORY it simply uses host fallback, which should usually be fine + but might have issues in corner cases.) + + We have default-device-var = 0 (default but also explicitly set), which will + fail at runtime. For -foffload=disable without non-host devices, see + target-50.c testcase. */ + +/* { dg-do run { target offload_device } } */ + +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "0" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +#include + +int +main () +{ + int x; + /* We know that there are non-host devices. With GCC, we still find them as + available devices, hence, check for it. */ + if (omp_get_num_devices() <= 0) + __builtin_abort (); + + /* But due to -foffload=disable, there are no binary code for (default) device '0' */ + + /* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot be used for offloading.*" } */ + /* { dg-shouldfail "OMP_TARGET_OFFLOAD=mandatory and no binary code for a non-host device" } */ + #pragma omp target map(tofrom:x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-51.c b/libgomp/testsuite/libgomp.c/target-51.c new file mode 100644 index 00000000000..7d09bceacd5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-51.c @@ -0,0 +1,24 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices, + which is enforced by using -foffload=disable. */ + +/* { dg-do run } */ +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ + +/* { dg-shouldfail "OMP_TARGET_OFFLOAD=mandatory and no available device" } */ + +/* See comment in target-50.c/target-50.c for why the output differs. */ + +/* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY but only the host device is available.*" { target { ! offload_device } } } */ +/* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY but device not found.*" { target offload_device } } */ + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-52.c b/libgomp/testsuite/libgomp.c/target-52.c new file mode 100644 index 00000000000..809380c6928 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-52.c @@ -0,0 +1,25 @@ +/* Only run this with available non-host devices; in that case, GCC sets + the default-device-var to 0. */ + +/* { dg-do run { target { offload_device } } } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +#include + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5 + omp_is_initial_device (); + + if (x != 5) + __builtin_abort (); + + if (0 != omp_get_default_device()) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-53.c b/libgomp/testsuite/libgomp.c/target-53.c new file mode 100644 index 00000000000..866e8961af1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-53.c @@ -0,0 +1,22 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "disabled" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '\[0-9\]+'.*OMP_TARGET_OFFLOAD = 'DISABLED'.*" } */ + +#include + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5 + omp_is_initial_device (); + + if (x != 5+1) + __builtin_abort (); + + if (omp_get_default_device() != omp_get_initial_device()) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-54.c b/libgomp/testsuite/libgomp.c/target-54.c new file mode 100644 index 00000000000..bc4e69b5278 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-54.c @@ -0,0 +1,20 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "default" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'DEFAULT'.*" } */ + +#include + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5 + omp_is_initial_device (); + + if (x != 5 + (omp_get_default_device() == omp_get_initial_device())) + __builtin_abort (); + + return 0; +}