| Message ID | 04045d90-625a-9214-31d4-0d00dbaf6ad7@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | openmp: Implement support for OMP_TARGET_OFFLOAD | expand |
On Wed, Oct 14, 2020 at 09:32:42PM +0100, Kwok Cheung Yeung wrote: > I'm not sure how this can be tested automatically, as the behaviour depends > on whether the compiler has been built with offloading support, and whether > any supported offloading hardware has been installed on the system. I have > not included any testcases for now. Not testing it is probably fine, unless we added some effective targets that the test would be run only if the offloading is (or is not present) and we would expect failures etc. I'm really not sure what the behavior should be in various cases though, so I've asked on omp-lang and will see what the committee members think. The fuzzy things are e.g.: void foo () {} #pragma omp declare target to (foo) int main () { #pragma omp target if(false) foo (); // Is this ok? omp_set_default_device (omp_get_initial_device ()); #pragma omp target foo (); // What about this? #pragma omp target device (omp_get_initial_device ()) foo (); // Or this? #pragma omp target device (omp_get_num_devices () + 42) foo (); // This one is clearly an error if (omp_get_num_devices () == 3) { #pragma omp target device (1) foo (); // This would be an error if we can't offload to device 1 } } Also, what about the case where there are no offloading devices at all, so the default device defaults to the initial device number? If nothing of this is valid (i.e. mandatory forbids all the fallback), then your patch implements it, but if e.g. only if (false) is allowed and nothing else, then we are ATM in trouble (we pass -2 for if false as well as it is returned from omp_get_initial_device (); thankfully, OpenMP 5.0 changed it (something that should be implemented in libgomp too), so that omp_get_initial_device () is now required to be omp_get_num_devices (), i.e. the host device is always the last one; so, we would do fallback for -2 and not for anything else), if also omp_get_initial_device () number is valid, then we should allow fallback for devices -2 and omp_get_initial_device () and punt on everything else. Jakub
On Thu, Oct 15, 2020 at 12:50:18PM +0200, Jakub Jelinek via Gcc-patches wrote: > If nothing of this is valid (i.e. mandatory forbids all the fallback), then > your patch implements it, but if e.g. only if (false) is allowed and nothing > else, then we are ATM in trouble (we pass -2 for if false as well as it is > returned from omp_get_initial_device (); thankfully, OpenMP 5.0 changed it > (something that should be implemented in libgomp too), so that > omp_get_initial_device () is now required to be omp_get_num_devices (), > i.e. the host device is always the last one; so, we would do fallback > for -2 and not for anything else), if also omp_get_initial_device () > number is valid, then we should allow fallback for devices -2 and > omp_get_initial_device () and punt on everything else. Ok, the first response indicates that both if (false) and omp_get_initial_device () are valid. Therefore, I think until omp_get_initial_device () value is changed, we want in resolve_device: if (device_id < 0 || device_id >= gomp_get_num_devices ()) { if (device_id != GOMP_DEVICE_HOST_FALLBACK) gomp_fatal (...); return NULL; } and do gomp_fatal also for further return NULLs in the function. And then in if (devicep == NULL || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) /* All shared memory devices should use the GOMP_target_ext function. */ || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM || !(fn_addr = gomp_get_target_fn_addr (devicep, fn))) return gomp_target_fallback (fn, hostaddrs); and similar do gomp_fatal before the gomp_target_fallback call if target-offload-var ICV is mandatory and devicep != NULL. And when we change omp_get_initial_device (), we'd use device_id > gomp_get_num_devices () above instead, and call gomp_fatal there only for (device_id != GOMP_DEVICE_HOST_FALLBACK && device_id != gomp_get_num_devices ()) Jakub
On 15/10/2020 12:02 pm, Jakub Jelinek wrote: > On Thu, Oct 15, 2020 at 12:50:18PM +0200, Jakub Jelinek via Gcc-patches wrote: > Ok, the first response indicates that both if (false) and > omp_get_initial_device () are valid. > Therefore, I think until omp_get_initial_device () value is changed, we > want in resolve_device: > if (device_id < 0 || device_id >= gomp_get_num_devices ()) > { > if (device_id != GOMP_DEVICE_HOST_FALLBACK) > gomp_fatal (...); > return NULL; > } > and do gomp_fatal also for further return NULLs in the function. > And then in > if (devicep == NULL > || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) > /* All shared memory devices should use the GOMP_target_ext function. */ > || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM > || !(fn_addr = gomp_get_target_fn_addr (devicep, fn))) > return gomp_target_fallback (fn, hostaddrs); > and similar do gomp_fatal before the gomp_target_fallback call > if target-offload-var ICV is mandatory and devicep != NULL. > Instead of doing a gomp_fatal before every call to gomp_target_fallback/gomp_target_data_fallback, I think it would be tidier to pass devicep into the fallback instead and keep the gomp_fatals in the fallback functions? Although it is kind of odd to pass the device to the fallback function to be used if the device in question does not work, but at least the fallback functions are static to target.c. Is this version okay for trunk? Thanks Kwok commit 82555f50d2930f973ab20782ebcb836b719bce96 Author: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Mon Oct 19 10:47:42 2020 -0700 openmp: Implement support for OMP_TARGET_OFFLOAD environment variable This implements support for the OMP_TARGET_OFFLOAD environment variable introduced in the OpenMP 5.0 standard, which controls how offloading is handled. It may be set to MANDATORY (abort if offloading cannot be performed), DISABLED (no offloading to devices) or DEFAULT (offload to device if possible, fall back to host if not). 2020-10-19 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_target_offload_var): New. (parse_target_offload): New. (handle_omp_display_env): Print value of OMP_TARGET_OFFLOAD. (initialize_env): Parse OMP_TARGET_OFFLOAD. * libgomp.h (gomp_target_offload_t): New. (gomp_target_offload_var): New. * libgomp.texi (OMP_TARGET_OFFLOAD): New section. * target.c (resolve_device): Generate error if device not found and offloading is mandatory. (gomp_target_fallback): Generate error if offloading is mandatory. (GOMP_target): Add argument in call to gomp_target_fallback. (GOMP_target_ext): Likewise. (gomp_target_data_fallback): Generate error if offloading is mandatory. (GOMP_target_data): Add argument in call to gomp_target_data_fallback. (GOMP_target_data_ext): Likewise. (gomp_target_task_fn): Add argument in call to gomp_target_fallback. (gomp_target_init): Return early if offloading is disabled. diff --git a/libgomp/env.c b/libgomp/env.c index d730c48..d0eae8d 100644 --- a/libgomp/env.c +++ b/libgomp/env.c @@ -75,6 +75,7 @@ struct gomp_task_icv gomp_global_icv = { unsigned long gomp_max_active_levels_var = gomp_supported_active_levels; bool gomp_cancel_var = false; +enum gomp_target_offload_t gomp_target_offload_var = GOMP_TARGET_OFFLOAD_DEFAULT; int gomp_max_task_priority_var = 0; #ifndef HAVE_SYNC_BUILTINS gomp_mutex_t gomp_managed_threads_lock; @@ -374,6 +375,48 @@ parse_unsigned_long_list (const char *name, unsigned long *p1stvalue, return false; } +static void +parse_target_offload (const char *name, enum gomp_target_offload_t *offload) +{ + const char *env; + bool found = false; + enum gomp_target_offload_t new_offload; + + env = getenv (name); + if (env == NULL) + return; + + while (isspace ((unsigned char) *env)) + ++env; + if (strncasecmp (env, "default", 7) == 0) + { + env += 7; + found = true; + new_offload = GOMP_TARGET_OFFLOAD_DEFAULT; + } + else if (strncasecmp (env, "mandatory", 9) == 0) + { + env += 9; + found = true; + new_offload = GOMP_TARGET_OFFLOAD_MANDATORY; + } + else if (strncasecmp (env, "disabled", 8) == 0) + { + env += 8; + found = true; + new_offload = GOMP_TARGET_OFFLOAD_DISABLED; + } + while (isspace ((unsigned char) *env)) + ++env; + if (found && *env == '\0') + { + *offload = new_offload; + return; + } + + gomp_error ("Invalid value for environment variable OMP_TARGET_OFFLOAD"); +} + /* Parse environment variable set to a boolean or list of omp_proc_bind_t enum values. Return true if one was present and it was successfully parsed. */ @@ -1334,6 +1377,21 @@ handle_omp_display_env (unsigned long stacksize, int wait_policy) } fputs ("'\n", stderr); + fputs (" OMP_TARGET_OFFLOAD = '", stderr); + switch (gomp_target_offload_var) + { + case GOMP_TARGET_OFFLOAD_DEFAULT: + fputs ("DEFAULT", stderr); + break; + case GOMP_TARGET_OFFLOAD_MANDATORY: + fputs ("MANDATORY", stderr); + break; + case GOMP_TARGET_OFFLOAD_DISABLED: + fputs ("DISABLED", stderr); + break; + } + fputs ("'\n", stderr); + if (verbose) { fputs (" GOMP_CPU_AFFINITY = ''\n", stderr); @@ -1366,6 +1424,7 @@ initialize_env (void) parse_boolean ("OMP_CANCELLATION", &gomp_cancel_var); parse_boolean ("OMP_DISPLAY_AFFINITY", &gomp_display_affinity_var); parse_int ("OMP_DEFAULT_DEVICE", &gomp_global_icv.default_device_var, true); + parse_target_offload ("OMP_TARGET_OFFLOAD", &gomp_target_offload_var); parse_int ("OMP_MAX_TASK_PRIORITY", &gomp_max_task_priority_var, true); parse_unsigned_long ("OMP_MAX_ACTIVE_LEVELS", &gomp_max_active_levels_var, true); diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 9d26de2..da7ac03 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -434,6 +434,13 @@ struct gomp_task_icv struct target_mem_desc *target_data; }; +enum gomp_target_offload_t +{ + GOMP_TARGET_OFFLOAD_DEFAULT, + GOMP_TARGET_OFFLOAD_MANDATORY, + GOMP_TARGET_OFFLOAD_DISABLED +}; + #define gomp_supported_active_levels INT_MAX extern struct gomp_task_icv gomp_global_icv; @@ -442,6 +449,7 @@ extern gomp_mutex_t gomp_managed_threads_lock; #endif extern unsigned long gomp_max_active_levels_var; extern bool gomp_cancel_var; +extern enum gomp_target_offload_t gomp_target_offload_var; extern int gomp_max_task_priority_var; extern unsigned long long gomp_spin_count_var, gomp_throttled_spin_count_var; extern unsigned long gomp_available_cpus, gomp_managed_threads; diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index a888613..7c6d5fd 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -1381,6 +1381,7 @@ beginning with @env{GOMP_} are GNU extensions. * OMP_PLACES:: Specifies on which CPUs the theads should be placed * OMP_STACKSIZE:: Set default thread stack size * OMP_SCHEDULE:: How threads are scheduled +* OMP_TARGET_OFFLOAD:: Controls offloading behaviour * OMP_THREAD_LIMIT:: Set the maximum number of threads * OMP_WAIT_POLICY:: How waiting threads are handled * GOMP_CPU_AFFINITY:: Bind threads to specific CPUs @@ -1654,6 +1655,30 @@ dynamic scheduling and a chunk size of 1 is used. +@node OMP_TARGET_OFFLOAD +@section @env{OMP_TARGET_OFFLOAD} -- Controls offloading behaviour +@cindex Environment Variable +@cindex Implementation specific setting +@table @asis +@item @emph{Description}: +Specifies the behaviour with regard to offloading code to a device. This +variable can be set to one of three values - @code{MANDATORY}, @code{DISABLED} +or @code{DEFAULT}. + +If set to @code{MANDATORY}, the program will terminate with an error if +the offload device is not present or is not supported. If set to +@code{DISABLED}, then offloading is disabled and all code will run on the +host. If set to @code{DEFAULT}, the program will try offloading to the +device first, then fall back to running code on the host if it cannot. + +If undefined, then the program will behave as if @code{DEFAULT} was set. + +@item @emph{Reference}: +@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 6.17 +@end table + + + @node OMP_THREAD_LIMIT @section @env{OMP_THREAD_LIMIT} -- Set the maximum number of threads @cindex Environment Variable diff --git a/libgomp/target.c b/libgomp/target.c index ab7ac9b..c568670 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -116,7 +116,13 @@ resolve_device (int device_id) } if (device_id < 0 || device_id >= gomp_get_num_devices ()) - return NULL; + { + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY + && device_id != GOMP_DEVICE_HOST_FALLBACK) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device not found."); + + return NULL; + } gomp_mutex_lock (&devices[device_id].lock); if (devices[device_id].state == GOMP_DEVICE_UNINITIALIZED) @@ -124,6 +130,11 @@ resolve_device (int device_id) else if (devices[device_id].state == GOMP_DEVICE_FINALIZED) { gomp_mutex_unlock (&devices[device_id].lock); + + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY + && device_id != GOMP_DEVICE_HOST_FALLBACK) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device is finalized."); + return NULL; } gomp_mutex_unlock (&devices[device_id].lock); @@ -1997,9 +2008,16 @@ gomp_unload_device (struct gomp_device_descr *devicep) /* Host fallback for GOMP_target{,_ext} routines. */ static void -gomp_target_fallback (void (*fn) (void *), void **hostaddrs) +gomp_target_fallback (void (*fn) (void *), void **hostaddrs, + struct gomp_device_descr *devicep) { struct gomp_thread old_thr, *thr = gomp_thread (); + + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY + && devicep != NULL) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot " + "be used for offloading."); + old_thr = *thr; memset (thr, '\0', sizeof (*thr)); if (gomp_places_list) @@ -2107,7 +2125,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused, /* All shared memory devices should use the GOMP_target_ext function. */ || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM || !(fn_addr = gomp_get_target_fn_addr (devicep, fn))) - return gomp_target_fallback (fn, hostaddrs); + return gomp_target_fallback (fn, hostaddrs, devicep); struct target_mem_desc *tgt_vars = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, @@ -2243,7 +2261,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, tgt_align, tgt_size); } } - gomp_target_fallback (fn, hostaddrs); + gomp_target_fallback (fn, hostaddrs, devicep); return; } @@ -2276,9 +2294,15 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, /* Host fallback for GOMP_target_data{,_ext} routines. */ static void -gomp_target_data_fallback (void) +gomp_target_data_fallback (struct gomp_device_descr *devicep) { struct gomp_task_icv *icv = gomp_icv (false); + + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY + && devicep != NULL) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot " + "be used for offloading."); + if (icv->target_data) { /* Even when doing a host fallback, if there are any active @@ -2302,7 +2326,7 @@ GOMP_target_data (int device, const void *unused, size_t mapnum, if (devicep == NULL || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) || (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)) - return gomp_target_data_fallback (); + return gomp_target_data_fallback (devicep); struct target_mem_desc *tgt = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, @@ -2321,7 +2345,7 @@ GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs, if (devicep == NULL || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) - return gomp_target_data_fallback (); + return gomp_target_data_fallback (devicep); struct target_mem_desc *tgt = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true, @@ -2617,7 +2641,7 @@ gomp_target_task_fn (void *data) || (devicep->can_run_func && !devicep->can_run_func (fn_addr))) { ttask->state = GOMP_TARGET_TASK_FALLBACK; - gomp_target_fallback (ttask->fn, ttask->hostaddrs); + gomp_target_fallback (ttask->fn, ttask->hostaddrs, devicep); return false; } @@ -3258,6 +3282,9 @@ gomp_target_init (void) num_devices = 0; devices = NULL; + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED) + return; + cur = OFFLOAD_PLUGINS; if (*cur) do
On Mon, Oct 19, 2020 at 06:57:49PM +0100, Kwok Cheung Yeung wrote: > --- a/libgomp/env.c > +++ b/libgomp/env.c > @@ -75,6 +75,7 @@ struct gomp_task_icv gomp_global_icv = { > > unsigned long gomp_max_active_levels_var = gomp_supported_active_levels; > bool gomp_cancel_var = false; > +enum gomp_target_offload_t gomp_target_offload_var = GOMP_TARGET_OFFLOAD_DEFAULT; Too long line. > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -116,7 +116,13 @@ resolve_device (int device_id) > } > > if (device_id < 0 || device_id >= gomp_get_num_devices ()) > - return NULL; > + { > + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY > + && device_id != GOMP_DEVICE_HOST_FALLBACK) > + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device not found."); No full stop at the end of gomp_fatal messages (everywhere in your patch) for consistency with other gomp_fatal calls. Also, too long line (similarly other gomp_fatal calls). Just use gomp_fatal (" ..... ... " ".........."); Otherwise LGTM. Jakub
On 10/19/20 8:21 PM, Jakub Jelinek via Gcc-patches wrote: > On Mon, Oct 19, 2020 at 06:57:49PM +0100, Kwok Cheung Yeung wrote: >> --- a/libgomp/target.c >> +++ b/libgomp/target.c ... > Otherwise LGTM. Unfortunately, the committed patch (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) causes build errors. The error seems to be provoked by function cloning – as the code itself looks fine: static void gomp_target_init (void) { ... num_devices = 0; ... // NEW LINES: if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED) return; ... struct gomp_device_descr *devices_s = malloc (num_devices * sizeof (struct gomp_device_descr)); ... for (i = 0; i < num_devices; i++) if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) devices_s[num_devices_after_openmp++] = devices[i]; I fail so see why this gives now the -Werror warning: ../../../repos/gcc/libgomp/target.c: In function ‘gomp_target_init.part.0’: ../../../repos/gcc/libgomp/target.c:3367:45: error: array subscript [-27450512014448737, 27450512014448737] is outside array bounds of ‘struct gomp_device_descr[0]’ [-Werror=array-bounds] 3367 | devices_s[num_devices_after_openmp++] = devices[i]; | ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~ ../../../repos/gcc/libgomp/target.c:3353:7: note: referencing an object of size 0 allocated by ‘malloc’ 3353 | = malloc (num_devices * sizeof (struct gomp_device_descr)); | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ../../../repos/gcc/libgomp/target.c:3363:39: error: array subscript [-27450512014448737, 27450512014448737] is outside array bounds of ‘struct gomp_device_descr[0]’ [-Werror=array-bounds] 3363 | devices_s[num_devices_openmp++] = devices[i]; | ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~ ../../../repos/gcc/libgomp/target.c:3353:7: note: referencing an object of size 0 allocated by ‘malloc’ 3353 | = malloc (num_devices * sizeof (struct gomp_device_descr)); | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On 10/20/20 2:11 PM, Tobias Burnus wrote: > Unfortunately, the committed patch > (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) > causes build errors. > > The error seems to be provoked by function cloning – as the code > itself looks fine: > ... > struct gomp_device_descr *devices_s > = malloc (num_devices * sizeof (struct gomp_device_descr)); > ... > for (i = 0; i < num_devices; i++) > if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) > devices_s[num_devices_after_openmp++] = devices[i]; gomp_target_init.part.0 () { ... <bb 2> devices_s_1 = malloc (0); ... num_devices.16_67 = num_devices; ... if (num_devices.16_67 > 0) goto <bb 3>; [89.00%] else goto <bb 18>; [11.00%] Which seems to have an ordering problem. Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On Tue, Oct 20, 2020 at 02:17:26PM +0200, Tobias Burnus wrote: > On 10/20/20 2:11 PM, Tobias Burnus wrote: > > > Unfortunately, the committed patch > > (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) > > causes build errors. > > > > The error seems to be provoked by function cloning – as the code > > itself looks fine: > > ... > > struct gomp_device_descr *devices_s > > = malloc (num_devices * sizeof (struct gomp_device_descr)); > > ... > > for (i = 0; i < num_devices; i++) > > if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) > > devices_s[num_devices_after_openmp++] = devices[i]; > > gomp_target_init.part.0 () > { > ... > <bb 2> > devices_s_1 = malloc (0); > ... > num_devices.16_67 = num_devices; > ... > if (num_devices.16_67 > 0) > goto <bb 3>; [89.00%] > else > goto <bb 18>; [11.00%] > > Which seems to have an ordering problem. This patch fixes the warning that breaks the bootstrap, but haven't tested it with offloading to see if it doesn't break offloading somehow. 2020-10-20 Jakub Jelinek <jakub@redhat.com> * target.c (gomp_target_init): Inside of the function, use automatic variables corresponding to num_devices, num_devices_openmp and devices global variables and update the globals only at the end of the function. --- libgomp/target.c.jj 2020-10-20 14:37:36.630967911 +0200 +++ libgomp/target.c 2020-10-20 14:52:36.556023803 +0200 @@ -3279,10 +3279,9 @@ gomp_target_init (void) const char *suffix = SONAME_SUFFIX (1); const char *cur, *next; char *plugin_name; - int i, new_num_devices; - - num_devices = 0; - devices = NULL; + int i, new_num_devs; + int num_devs = 0, num_devs_openmp; + struct gomp_device_descr *devs = NULL; if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED) return; @@ -3303,7 +3302,7 @@ gomp_target_init (void) plugin_name = (char *) malloc (prefix_len + cur_len + suffix_len + 1); if (!plugin_name) { - num_devices = 0; + num_devs = 0; break; } @@ -3313,16 +3312,16 @@ gomp_target_init (void) if (gomp_load_plugin_for_device (¤t_device, plugin_name)) { - new_num_devices = current_device.get_num_devices_func (); - if (new_num_devices >= 1) + new_num_devs = current_device.get_num_devices_func (); + if (new_num_devs >= 1) { /* Augment DEVICES and NUM_DEVICES. */ - devices = realloc (devices, (num_devices + new_num_devices) - * sizeof (struct gomp_device_descr)); - if (!devices) + devs = realloc (devs, (num_devs + new_num_devs) + * sizeof (struct gomp_device_descr)); + if (!devs) { - num_devices = 0; + num_devs = 0; free (plugin_name); break; } @@ -3332,12 +3331,12 @@ gomp_target_init (void) current_device.type = current_device.get_type_func (); current_device.mem_map.root = NULL; current_device.state = GOMP_DEVICE_UNINITIALIZED; - for (i = 0; i < new_num_devices; i++) + for (i = 0; i < new_num_devs; i++) { current_device.target_id = i; - devices[num_devices] = current_device; - gomp_mutex_init (&devices[num_devices].lock); - num_devices++; + devs[num_devs] = current_device; + gomp_mutex_init (&devs[num_devs].lock); + num_devs++; } } } @@ -3349,34 +3348,37 @@ gomp_target_init (void) /* In DEVICES, sort the GOMP_OFFLOAD_CAP_OPENMP_400 ones first, and set NUM_DEVICES_OPENMP. */ - struct gomp_device_descr *devices_s - = malloc (num_devices * sizeof (struct gomp_device_descr)); - if (!devices_s) + struct gomp_device_descr *devs_s + = malloc (num_devs * sizeof (struct gomp_device_descr)); + if (!devs_s) { - num_devices = 0; - free (devices); - devices = NULL; + num_devs = 0; + free (devs); + devs = NULL; } - num_devices_openmp = 0; - for (i = 0; i < num_devices; i++) - if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) - devices_s[num_devices_openmp++] = devices[i]; - int num_devices_after_openmp = num_devices_openmp; - for (i = 0; i < num_devices; i++) - if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) - devices_s[num_devices_after_openmp++] = devices[i]; - free (devices); - devices = devices_s; + num_devs_openmp = 0; + for (i = 0; i < num_devs; i++) + if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) + devs_s[num_devs_openmp++] = devs[i]; + int num_devs_after_openmp = num_devs_openmp; + for (i = 0; i < num_devs; i++) + if (!(devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) + devs_s[num_devs_after_openmp++] = devs[i]; + free (devs); + devs = devs_s; - for (i = 0; i < num_devices; i++) + for (i = 0; i < num_devs; i++) { /* The 'devices' array can be moved (by the realloc call) until we have found all the plugins, so registering with the OpenACC runtime (which takes a copy of the pointer argument) must be delayed until now. */ - if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200) - goacc_register (&devices[i]); + if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200) + goacc_register (&devs[i]); } + num_devices = num_devs; + num_devices_openmp = num_devs_openmp; + devices = devs; if (atexit (gomp_target_fini) != 0) gomp_fatal ("atexit failed"); } Jakub
Hi Tobias, > On 10/19/20 8:21 PM, Jakub Jelinek via Gcc-patches wrote: > >> On Mon, Oct 19, 2020 at 06:57:49PM +0100, Kwok Cheung Yeung wrote: >>> --- a/libgomp/target.c >>> +++ b/libgomp/target.c > ... >> Otherwise LGTM. > > Unfortunately, the committed patch (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) > causes build errors. the patch also breaks bootstrap on both i386-pc-solaris2.11 and sparc-sun-solaris2.11: /vol/gcc/src/hg/master/local/libgomp/env.c: In function 'initialize_env': /vol/gcc/src/hg/master/local/libgomp/env.c:414:16: error: 'new_offload' may be used uninitialized in this function [-Werror=maybe-uninitialized] 414 | *offload = new_offload; | ~~~~~~~~~^~~~~~~~~~~~~ /vol/gcc/src/hg/master/local/libgomp/env.c:384:30: note: 'new_offload' was declared here 384 | enum gomp_target_offload_t new_offload; | ^~~~~~~~~~~ Rainer
On Tue, Oct 20, 2020 at 06:39:39PM +0200, Rainer Orth wrote: > Hi Tobias, > > > On 10/19/20 8:21 PM, Jakub Jelinek via Gcc-patches wrote: > > > >> On Mon, Oct 19, 2020 at 06:57:49PM +0100, Kwok Cheung Yeung wrote: > >>> --- a/libgomp/target.c > >>> +++ b/libgomp/target.c > > ... > >> Otherwise LGTM. > > > > Unfortunately, the committed patch (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) > > causes build errors. > > the patch also breaks bootstrap on both i386-pc-solaris2.11 and > sparc-sun-solaris2.11: > > /vol/gcc/src/hg/master/local/libgomp/env.c: In function 'initialize_env': > /vol/gcc/src/hg/master/local/libgomp/env.c:414:16: error: 'new_offload' may be used uninitialized in this function [-Werror=maybe-uninitialized] > 414 | *offload = new_offload; > | ~~~~~~~~~^~~~~~~~~~~~~ > /vol/gcc/src/hg/master/local/libgomp/env.c:384:30: note: 'new_offload' was declared here > 384 | enum gomp_target_offload_t new_offload; > | ^~~~~~~~~~~ I can't reproduce that, but I fail to see why we need two separate variables, one with actual value and one tracking if the value is valid. So I'd go with: 2020-10-20 Jakub Jelinek <jakub@redhat.com> * env.c (parse_target_offload): Change new_offload var type to int, preinitialize to -1, remove found var and test new_offload != -1 instead of found. --- libgomp/env.c.jj 2020-10-20 14:37:36.593968443 +0200 +++ libgomp/env.c 2020-10-20 18:43:00.338389023 +0200 @@ -380,8 +380,7 @@ static void parse_target_offload (const char *name, enum gomp_target_offload_t *offload) { const char *env; - bool found = false; - enum gomp_target_offload_t new_offload; + int new_offload = -1; env = getenv (name); if (env == NULL) @@ -392,24 +391,21 @@ parse_target_offload (const char *name, if (strncasecmp (env, "default", 7) == 0) { env += 7; - found = true; new_offload = GOMP_TARGET_OFFLOAD_DEFAULT; } else if (strncasecmp (env, "mandatory", 9) == 0) { env += 9; - found = true; new_offload = GOMP_TARGET_OFFLOAD_MANDATORY; } else if (strncasecmp (env, "disabled", 8) == 0) { env += 8; - found = true; new_offload = GOMP_TARGET_OFFLOAD_DISABLED; } while (isspace ((unsigned char) *env)) ++env; - if (found && *env == '\0') + if (new_offload != -1 && *env == '\0') { *offload = new_offload; return; Jakub
On 20/10/2020 1:57 pm, Jakub Jelinek wrote: > On Tue, Oct 20, 2020 at 02:17:26PM +0200, Tobias Burnus wrote: >> On 10/20/20 2:11 PM, Tobias Burnus wrote: >> >>> Unfortunately, the committed patch >>> (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) >>> causes build errors. >>> >>> The error seems to be provoked by function cloning – as the code >>> itself looks fine: >>> ... >>> struct gomp_device_descr *devices_s >>> = malloc (num_devices * sizeof (struct gomp_device_descr)); >>> ... >>> for (i = 0; i < num_devices; i++) >>> if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) >>> devices_s[num_devices_after_openmp++] = devices[i]; >> >> gomp_target_init.part.0 () >> { >> ... >> <bb 2> >> devices_s_1 = malloc (0); >> ... >> num_devices.16_67 = num_devices; >> ... >> if (num_devices.16_67 > 0) >> goto <bb 3>; [89.00%] >> else >> goto <bb 18>; [11.00%] >> >> Which seems to have an ordering problem. > > This patch fixes the warning that breaks the bootstrap, but haven't > tested it with offloading to see if it doesn't break offloading somehow. > Thank you for the fix. It appears that the issue only occurs when offloading is not enabled (I tested with offloading to Nvidia and AMD GCN devices, but forgot to test no offloading). Kwok
Hi Jakub, >> the patch also breaks bootstrap on both i386-pc-solaris2.11 and >> sparc-sun-solaris2.11: >> >> /vol/gcc/src/hg/master/local/libgomp/env.c: In function 'initialize_env': >> /vol/gcc/src/hg/master/local/libgomp/env.c:414:16: error: 'new_offload' >> may be used uninitialized in this function [-Werror=maybe-uninitialized] >> 414 | *offload = new_offload; >> | ~~~~~~~~~^~~~~~~~~~~~~ >> /vol/gcc/src/hg/master/local/libgomp/env.c:384:30: note: 'new_offload' >> was declared here >> 384 | enum gomp_target_offload_t new_offload; >> | ^~~~~~~~~~~ > > I can't reproduce that, but I fail to see why we need two separate > variables, one with actual value and one tracking if the value is valid. I'd also tried i686-pc-linux-gnu, but didn't get the failure there either. > So I'd go with: > > 2020-10-20 Jakub Jelinek <jakub@redhat.com> > > * env.c (parse_target_offload): Change new_offload var type to int, > preinitialize to -1, remove found var and test new_offload != -1 > instead of found. That worked just fine. Thanks. Rainer
diff --git a/libgomp/env.c b/libgomp/env.c index d730c48..d0eae8d 100644 --- a/libgomp/env.c +++ b/libgomp/env.c @@ -75,6 +75,7 @@ struct gomp_task_icv gomp_global_icv = { unsigned long gomp_max_active_levels_var = gomp_supported_active_levels; bool gomp_cancel_var = false; +enum gomp_target_offload_t gomp_target_offload_var = GOMP_TARGET_OFFLOAD_DEFAULT; int gomp_max_task_priority_var = 0; #ifndef HAVE_SYNC_BUILTINS gomp_mutex_t gomp_managed_threads_lock; @@ -374,6 +375,48 @@ parse_unsigned_long_list (const char *name, unsigned long *p1stvalue, return false; } +static void +parse_target_offload (const char *name, enum gomp_target_offload_t *offload) +{ + const char *env; + bool found = false; + enum gomp_target_offload_t new_offload; + + env = getenv (name); + if (env == NULL) + return; + + while (isspace ((unsigned char) *env)) + ++env; + if (strncasecmp (env, "default", 7) == 0) + { + env += 7; + found = true; + new_offload = GOMP_TARGET_OFFLOAD_DEFAULT; + } + else if (strncasecmp (env, "mandatory", 9) == 0) + { + env += 9; + found = true; + new_offload = GOMP_TARGET_OFFLOAD_MANDATORY; + } + else if (strncasecmp (env, "disabled", 8) == 0) + { + env += 8; + found = true; + new_offload = GOMP_TARGET_OFFLOAD_DISABLED; + } + while (isspace ((unsigned char) *env)) + ++env; + if (found && *env == '\0') + { + *offload = new_offload; + return; + } + + gomp_error ("Invalid value for environment variable OMP_TARGET_OFFLOAD"); +} + /* Parse environment variable set to a boolean or list of omp_proc_bind_t enum values. Return true if one was present and it was successfully parsed. */ @@ -1334,6 +1377,21 @@ handle_omp_display_env (unsigned long stacksize, int wait_policy) } fputs ("'\n", stderr); + fputs (" OMP_TARGET_OFFLOAD = '", stderr); + switch (gomp_target_offload_var) + { + case GOMP_TARGET_OFFLOAD_DEFAULT: + fputs ("DEFAULT", stderr); + break; + case GOMP_TARGET_OFFLOAD_MANDATORY: + fputs ("MANDATORY", stderr); + break; + case GOMP_TARGET_OFFLOAD_DISABLED: + fputs ("DISABLED", stderr); + break; + } + fputs ("'\n", stderr); + if (verbose) { fputs (" GOMP_CPU_AFFINITY = ''\n", stderr); @@ -1366,6 +1424,7 @@ initialize_env (void) parse_boolean ("OMP_CANCELLATION", &gomp_cancel_var); parse_boolean ("OMP_DISPLAY_AFFINITY", &gomp_display_affinity_var); parse_int ("OMP_DEFAULT_DEVICE", &gomp_global_icv.default_device_var, true); + parse_target_offload ("OMP_TARGET_OFFLOAD", &gomp_target_offload_var); parse_int ("OMP_MAX_TASK_PRIORITY", &gomp_max_task_priority_var, true); parse_unsigned_long ("OMP_MAX_ACTIVE_LEVELS", &gomp_max_active_levels_var, true); diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 9d26de2..da7ac03 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -434,6 +434,13 @@ struct gomp_task_icv struct target_mem_desc *target_data; }; +enum gomp_target_offload_t +{ + GOMP_TARGET_OFFLOAD_DEFAULT, + GOMP_TARGET_OFFLOAD_MANDATORY, + GOMP_TARGET_OFFLOAD_DISABLED +}; + #define gomp_supported_active_levels INT_MAX extern struct gomp_task_icv gomp_global_icv; @@ -442,6 +449,7 @@ extern gomp_mutex_t gomp_managed_threads_lock; #endif extern unsigned long gomp_max_active_levels_var; extern bool gomp_cancel_var; +extern enum gomp_target_offload_t gomp_target_offload_var; extern int gomp_max_task_priority_var; extern unsigned long long gomp_spin_count_var, gomp_throttled_spin_count_var; extern unsigned long gomp_available_cpus, gomp_managed_threads; diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 1c34bbe..4a5e56f 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -1381,6 +1381,7 @@ beginning with @env{GOMP_} are GNU extensions. * OMP_PLACES:: Specifies on which CPUs the theads should be placed * OMP_STACKSIZE:: Set default thread stack size * OMP_SCHEDULE:: How threads are scheduled +* OMP_TARGET_OFFLOAD:: Controls offloading behaviour * OMP_THREAD_LIMIT:: Set the maximum number of threads * OMP_WAIT_POLICY:: How waiting threads are handled * GOMP_CPU_AFFINITY:: Bind threads to specific CPUs @@ -1654,6 +1655,30 @@ dynamic scheduling and a chunk size of 1 is used. +@node OMP_TARGET_OFFLOAD +@section @env{OMP_TARGET_OFFLOAD} -- Controls offloading behaviour +@cindex Environment Variable +@cindex Implementation specific setting +@table @asis +@item @emph{Description}: +Specifies the behaviour with regard to offloading code to a device. This +variable can be set to one of three values - @code{MANDATORY}, @code{DISABLED} +or @code{DEFAULT}. + +If set to @code{MANDATORY}, the program will terminate with an error if +the offload device is not present or is not supported. If set to +@code{DISABLED}, then offloading is disabled and all code will run on the +host. If set to @code{DEFAULT}, the program will try offloading to the +device first, then fall back to running code on the host if it cannot. + +If undefined, then the program will behave as if @code{DEFAULT} was set. + +@item @emph{Reference}: +@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 6.17 +@end table + + + @node OMP_THREAD_LIMIT @section @env{OMP_THREAD_LIMIT} -- Set the maximum number of threads @cindex Environment Variable diff --git a/libgomp/target.c b/libgomp/target.c index ab7ac9b..8fcd7f8 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -116,7 +116,12 @@ resolve_device (int device_id) } if (device_id < 0 || device_id >= gomp_get_num_devices ()) - return NULL; + { + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device not found."); + + return NULL; + } gomp_mutex_lock (&devices[device_id].lock); if (devices[device_id].state == GOMP_DEVICE_UNINITIALIZED) @@ -2000,6 +2005,11 @@ static void gomp_target_fallback (void (*fn) (void *), void **hostaddrs) { struct gomp_thread old_thr, *thr = gomp_thread (); + + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot " + "be used for offloading."); + old_thr = *thr; memset (thr, '\0', sizeof (*thr)); if (gomp_places_list) @@ -2279,6 +2289,11 @@ static void gomp_target_data_fallback (void) { struct gomp_task_icv *icv = gomp_icv (false); + + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot " + "be used for offloading."); + if (icv->target_data) { /* Even when doing a host fallback, if there are any active @@ -3258,6 +3273,9 @@ gomp_target_init (void) num_devices = 0; devices = NULL; + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED) + return; + cur = OFFLOAD_PLUGINS; if (*cur) do
Hello This implements support for the OMP_TARGET_OFFLOAD environment variable introduced in the OpenMP 5.0 standard, which controls how offloading is handled in an OpenMP program. If set to MANDATORY, then libgomp will cause the program to abort with a gomp_fatal if an offload device is not found, or if it falls back to the host for some reason. When DISABLED, then gomp_target_init will return early, so that libgomp acts as if no offload devices were found and the host fallback is always used. For DEFAULT, nothing is done, resulting in the original behaviour. I'm not sure how this can be tested automatically, as the behaviour depends on whether the compiler has been built with offloading support, and whether any supported offloading hardware has been installed on the system. I have not included any testcases for now. Okay for trunk? Thanks Kwok commit a22f434d5ec9e62c158912b693275ce89a2cbab0 Author: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Thu Oct 8 10:08:27 2020 -0700 openmp: Implement support for OMP_TARGET_OFFLOAD environment variable This implements support for the OMP_TARGET_OFFLOAD environment variable introduced in the OpenMP 5.0 standard, which controls how offloading is handled. It may be set to MANDATORY (abort if offloading cannot be performed), DISABLED (no offloading to devices) or DEFAULT (offload to device if possible, fall back to host if not). 2020-10-14 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_target_offload_var): New. (parse_target_offload): New. (handle_omp_display_env): Print value of OMP_TARGET_OFFLOAD. (initialize_env): Parse OMP_TARGET_OFFLOAD. * libgomp.h (gomp_target_offload_t): New. (gomp_target_offload_var): New. * libgomp.texi (OMP_TARGET_OFFLOAD): New section. * target.c (resolve_device): Generate error if device not found and offloading is mandatory. (gomp_target_fallback): Generate error if offloading is mandatory. (gomp_target_fallback): Likewise. (gomp_target_init): Return early if offloading is disabled.