* [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
@ 2020-10-14 20:32 Kwok Cheung Yeung
2020-10-15 10:50 ` Jakub Jelinek
0 siblings, 1 reply; 13+ messages in thread
From: Kwok Cheung Yeung @ 2020-10-14 20:32 UTC (permalink / raw)
To: GCC Patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 873 bytes --]
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
[-- Attachment #2: omp_target_offload.patch --]
[-- Type: text/plain, Size: 8290 bytes --]
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.
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-14 20:32 [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD Kwok Cheung Yeung
@ 2020-10-15 10:50 ` Jakub Jelinek
2020-10-15 11:02 ` Jakub Jelinek
0 siblings, 1 reply; 13+ messages in thread
From: Jakub Jelinek @ 2020-10-15 10:50 UTC (permalink / raw)
To: Kwok Cheung Yeung; +Cc: GCC Patches
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-15 10:50 ` Jakub Jelinek
@ 2020-10-15 11:02 ` Jakub Jelinek
2020-10-19 17:57 ` Kwok Cheung Yeung
2020-10-21 11:52 ` [PATCH] openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirements Jakub Jelinek
0 siblings, 2 replies; 13+ messages in thread
From: Jakub Jelinek @ 2020-10-15 11:02 UTC (permalink / raw)
To: Kwok Cheung Yeung, GCC Patches
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-15 11:02 ` Jakub Jelinek
@ 2020-10-19 17:57 ` Kwok Cheung Yeung
2020-10-19 18:21 ` Jakub Jelinek
2020-10-21 11:52 ` [PATCH] openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirements Jakub Jelinek
1 sibling, 1 reply; 13+ messages in thread
From: Kwok Cheung Yeung @ 2020-10-19 17:57 UTC (permalink / raw)
To: Jakub Jelinek, GCC Patches
[-- Attachment #1: Type: text/plain, Size: 1525 bytes --]
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
[-- Attachment #2: omp_target_offload.patch --]
[-- Type: text/plain, Size: 11550 bytes --]
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-19 17:57 ` Kwok Cheung Yeung
@ 2020-10-19 18:21 ` Jakub Jelinek
2020-10-20 12:11 ` Tobias Burnus
0 siblings, 1 reply; 13+ messages in thread
From: Jakub Jelinek @ 2020-10-19 18:21 UTC (permalink / raw)
To: Kwok Cheung Yeung; +Cc: GCC Patches
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-19 18:21 ` Jakub Jelinek
@ 2020-10-20 12:11 ` Tobias Burnus
2020-10-20 12:17 ` Tobias Burnus
2020-10-20 16:39 ` Rainer Orth
0 siblings, 2 replies; 13+ messages in thread
From: Tobias Burnus @ 2020-10-20 12:11 UTC (permalink / raw)
To: Jakub Jelinek, Kwok Cheung Yeung; +Cc: GCC Patches
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-20 12:11 ` Tobias Burnus
@ 2020-10-20 12:17 ` Tobias Burnus
2020-10-20 12:57 ` Jakub Jelinek
2020-10-20 16:39 ` Rainer Orth
1 sibling, 1 reply; 13+ messages in thread
From: Tobias Burnus @ 2020-10-20 12:17 UTC (permalink / raw)
To: Jakub Jelinek, Martin Jambor, Jan Hubicka; +Cc: GCC Patches, Kwok Cheung Yeung
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-20 12:17 ` Tobias Burnus
@ 2020-10-20 12:57 ` Jakub Jelinek
2020-10-20 21:28 ` Kwok Cheung Yeung
0 siblings, 1 reply; 13+ messages in thread
From: Jakub Jelinek @ 2020-10-20 12:57 UTC (permalink / raw)
To: Tobias Burnus; +Cc: Martin Jambor, Jan Hubicka, GCC Patches, Kwok Cheung Yeung
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-20 12:11 ` Tobias Burnus
2020-10-20 12:17 ` Tobias Burnus
@ 2020-10-20 16:39 ` Rainer Orth
2020-10-20 16:45 ` Jakub Jelinek
1 sibling, 1 reply; 13+ messages in thread
From: Rainer Orth @ 2020-10-20 16:39 UTC (permalink / raw)
To: Tobias Burnus; +Cc: Jakub Jelinek, Kwok Cheung Yeung, GCC Patches
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
--
-----------------------------------------------------------------------------
Rainer Orth, Center for Biotechnology, Bielefeld University
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-20 16:39 ` Rainer Orth
@ 2020-10-20 16:45 ` Jakub Jelinek
2020-10-22 10:44 ` Rainer Orth
0 siblings, 1 reply; 13+ messages in thread
From: Jakub Jelinek @ 2020-10-20 16:45 UTC (permalink / raw)
To: Rainer Orth; +Cc: Tobias Burnus, Kwok Cheung Yeung, GCC Patches
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-20 12:57 ` Jakub Jelinek
@ 2020-10-20 21:28 ` Kwok Cheung Yeung
0 siblings, 0 replies; 13+ messages in thread
From: Kwok Cheung Yeung @ 2020-10-20 21:28 UTC (permalink / raw)
To: Jakub Jelinek, Tobias Burnus; +Cc: Martin Jambor, Jan Hubicka, GCC Patches
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
^ permalink raw reply [flat|nested] 13+ messages in thread
* [PATCH] openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirements
2020-10-15 11:02 ` Jakub Jelinek
2020-10-19 17:57 ` Kwok Cheung Yeung
@ 2020-10-21 11:52 ` Jakub Jelinek
1 sibling, 0 replies; 13+ messages in thread
From: Jakub Jelinek @ 2020-10-21 11:52 UTC (permalink / raw)
To: Tobias Burnus, Kwok Cheung Yeung; +Cc: gcc-patches
On Thu, Oct 15, 2020 at 01:02:29PM +0200, Jakub Jelinek via Gcc-patches wrote:
> Therefore, I think until omp_get_initial_device () value is changed, we
The following so far untested patch implements that change.
OpenMP 4.5 said for omp_get_initial_device:
The value of the device number is implementation defined. If it is between 0 and one less than
omp_get_num_devices() then it is valid for use with all device constructs and routines; if it is
outside that range, then it is only valid for use with the device memory routines and not in the
device clause.
and OpenMP 5.0 similarly, but OpenMP 5.1 says:
The value of the device number is the value returned by the omp_get_num_devices routine.
As the new value is compatible with what has been required earlier, I think
we can change it already now.
2020-10-21 Jakub Jelinek <jakub@redhat.com>
* icv.c (omp_get_initial_device): Remove including corresponding
ialias.
* icv-device.c (omp_get_initial_device): New function. Return
gomp_get_num_devices (). Add ialias.
* target.c (resolve_device): Don't fail with
OMP_TARGET_OFFLOAD=mandatory if device_id is equal to
gomp_get_num_devices ().
(omp_target_alloc, omp_target_free, omp_target_is_present,
omp_target_memcpy, omp_target_memcpy_rect, omp_target_associate_ptr,
omp_target_disassociate_ptr, omp_pause_resource): Use
gomp_get_num_devices () instead of GOMP_DEVICE_HOST_FALLBACK on the
first use in the functions, in uses dominated by the
gomp_get_num_devices call use num_devices_openmp instead.
* libgomp.texi (omp_get_initial_device): Document.
* config/gcn/icv-device.c (omp_get_initial_device): New function.
Add ialias.
* config/nvptx/icv-device.c (omp_get_initial_device): Likewise.
* testsuite/libgomp.c/target-40.c: New test.
--- libgomp/icv.c.jj 2020-10-13 22:29:22.202958364 +0200
+++ libgomp/icv.c 2020-10-21 12:31:05.941289010 +0200
@@ -156,12 +156,6 @@ omp_get_proc_bind (void)
}
int
-omp_get_initial_device (void)
-{
- return GOMP_DEVICE_HOST_FALLBACK;
-}
-
-int
omp_get_num_places (void)
{
return gomp_places_list_len;
@@ -241,7 +235,6 @@ ialias (omp_get_max_active_levels)
ialias (omp_get_supported_active_levels)
ialias (omp_get_cancellation)
ialias (omp_get_proc_bind)
-ialias (omp_get_initial_device)
ialias (omp_get_max_task_priority)
ialias (omp_get_num_places)
ialias (omp_get_place_num)
--- libgomp/icv-device.c.jj 2020-01-12 11:54:39.016374137 +0100
+++ libgomp/icv-device.c 2020-10-21 12:32:24.827131320 +0200
@@ -43,6 +43,12 @@ omp_get_default_device (void)
}
int
+omp_get_initial_device (void)
+{
+ return gomp_get_num_devices ();
+}
+
+int
omp_get_num_devices (void)
{
return gomp_get_num_devices ();
@@ -57,5 +63,6 @@ omp_is_initial_device (void)
ialias (omp_set_default_device)
ialias (omp_get_default_device)
+ialias (omp_get_initial_device)
ialias (omp_get_num_devices)
ialias (omp_is_initial_device)
--- libgomp/target.c.jj 2020-10-20 19:51:38.149361531 +0200
+++ libgomp/target.c 2020-10-21 12:43:19.336526122 +0200
@@ -118,7 +118,8 @@ resolve_device (int device_id)
if (device_id < 0 || device_id >= gomp_get_num_devices ())
{
if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
- && device_id != GOMP_DEVICE_HOST_FALLBACK)
+ && device_id != GOMP_DEVICE_HOST_FALLBACK
+ && device_id != num_devices_openmp)
gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
"but device not found");
@@ -132,8 +133,7 @@ resolve_device (int device_id)
{
gomp_mutex_unlock (&devices[device_id].lock);
- if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
- && device_id != GOMP_DEVICE_HOST_FALLBACK)
+ if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY)
gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
"but device is finalized");
@@ -2716,7 +2716,7 @@ GOMP_teams (unsigned int num_teams, unsi
void *
omp_target_alloc (size_t size, int device_num)
{
- if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+ if (device_num == gomp_get_num_devices ())
return malloc (size);
if (device_num < 0)
@@ -2742,7 +2742,7 @@ omp_target_free (void *device_ptr, int d
if (device_ptr == NULL)
return;
- if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+ if (device_num == gomp_get_num_devices ())
{
free (device_ptr);
return;
@@ -2773,7 +2773,7 @@ omp_target_is_present (const void *ptr,
if (ptr == NULL)
return 1;
- if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+ if (device_num == gomp_get_num_devices ())
return 1;
if (device_num < 0)
@@ -2807,7 +2807,7 @@ omp_target_memcpy (void *dst, const void
struct gomp_device_descr *dst_devicep = NULL, *src_devicep = NULL;
bool ret;
- if (dst_device_num != GOMP_DEVICE_HOST_FALLBACK)
+ if (dst_device_num != gomp_get_num_devices ())
{
if (dst_device_num < 0)
return EINVAL;
@@ -2820,7 +2820,7 @@ omp_target_memcpy (void *dst, const void
|| dst_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
dst_devicep = NULL;
}
- if (src_device_num != GOMP_DEVICE_HOST_FALLBACK)
+ if (src_device_num != num_devices_openmp)
{
if (src_device_num < 0)
return EINVAL;
@@ -2958,7 +2958,7 @@ omp_target_memcpy_rect (void *dst, const
if (!dst && !src)
return INT_MAX;
- if (dst_device_num != GOMP_DEVICE_HOST_FALLBACK)
+ if (dst_device_num != gomp_get_num_devices ())
{
if (dst_device_num < 0)
return EINVAL;
@@ -2971,7 +2971,7 @@ omp_target_memcpy_rect (void *dst, const
|| dst_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
dst_devicep = NULL;
}
- if (src_device_num != GOMP_DEVICE_HOST_FALLBACK)
+ if (src_device_num != num_devices_openmp)
{
if (src_device_num < 0)
return EINVAL;
@@ -3007,7 +3007,7 @@ 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_DEVICE_HOST_FALLBACK)
+ if (device_num == gomp_get_num_devices ())
return EINVAL;
if (device_num < 0)
@@ -3070,7 +3070,7 @@ omp_target_associate_ptr (const void *ho
int
omp_target_disassociate_ptr (const void *ptr, int device_num)
{
- if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+ if (device_num == gomp_get_num_devices ())
return EINVAL;
if (device_num < 0)
@@ -3113,9 +3113,9 @@ int
omp_pause_resource (omp_pause_resource_t kind, int device_num)
{
(void) kind;
- if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+ if (device_num == gomp_get_num_devices ())
return gomp_pause_host ();
- if (device_num < 0 || device_num >= gomp_get_num_devices ())
+ if (device_num < 0 || device_num >= num_devices_openmp)
return -1;
/* Do nothing for target devices for now. */
return 0;
--- libgomp/libgomp.texi.jj 2020-10-20 14:37:36.594968429 +0200
+++ libgomp/libgomp.texi 2020-10-21 13:06:31.593116640 +0200
@@ -166,6 +166,7 @@ linkage, and do not throw exceptions.
* omp_get_cancellation:: Whether cancellation support is enabled
* omp_get_default_device:: Get the default device for target regions
* omp_get_dynamic:: Dynamic teams setting
+* omp_get_initial_device:: Device number of host device
* omp_get_level:: Number of parallel regions
* omp_get_max_active_levels:: Current maximum number of active regions
* omp_get_max_task_priority:: Maximum task priority value that can be set
@@ -352,6 +353,33 @@ disabled by default.
@end table
+
+@node omp_get_initial_device
+@section @code{omp_get_initial_device} -- Return device number of initial device
+@table @asis
+@item @emph{Description}:
+This function returns a device number that represents the host device.
+For OpenMP 5.1, this must be equal to the value returned by the
+@code{omp_get_num_devices} function.
+
+@item @emph{C/C++}
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_initial_device(void);}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_initial_device()}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_num_devices}
+
+@item @emph{Reference}:
+@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.35.
+@end table
+
+
@node omp_get_level
@section @code{omp_get_level} -- Obtain the current nesting level
--- libgomp/config/gcn/icv-device.c.jj 2020-01-12 11:54:39.012374197 +0100
+++ libgomp/config/gcn/icv-device.c 2020-10-21 12:32:54.423696983 +0200
@@ -40,6 +40,12 @@ omp_get_default_device (void)
}
int
+omp_get_initial_device (void)
+{
+ return GOMP_DEVICE_HOST_FALLBACK;
+}
+
+int
omp_get_num_devices (void)
{
return 0;
@@ -66,6 +72,7 @@ omp_is_initial_device (void)
ialias (omp_set_default_device)
ialias (omp_get_default_device)
+ialias (omp_get_initial_device)
ialias (omp_get_num_devices)
ialias (omp_get_num_teams)
ialias (omp_get_team_num)
--- libgomp/config/nvptx/icv-device.c.jj 2020-01-12 11:54:39.013374182 +0100
+++ libgomp/config/nvptx/icv-device.c 2020-10-21 12:33:29.757178445 +0200
@@ -40,6 +40,12 @@ omp_get_default_device (void)
}
int
+omp_get_initial_device (void)
+{
+ return GOMP_DEVICE_HOST_FALLBACK;
+}
+
+int
omp_get_num_devices (void)
{
return 0;
@@ -54,5 +60,6 @@ omp_is_initial_device (void)
ialias (omp_set_default_device)
ialias (omp_get_default_device)
+ialias (omp_get_initial_device)
ialias (omp_get_num_devices)
ialias (omp_is_initial_device)
--- libgomp/testsuite/libgomp.c/target-40.c.jj 2020-10-21 13:10:55.658250523 +0200
+++ libgomp/testsuite/libgomp.c/target-40.c 2020-10-21 13:10:49.348342902 +0200
@@ -0,0 +1,10 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ if (omp_get_initial_device () != omp_get_num_devices ())
+ abort ();
+ return 0;
+}
Jakub
^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD
2020-10-20 16:45 ` Jakub Jelinek
@ 2020-10-22 10:44 ` Rainer Orth
0 siblings, 0 replies; 13+ messages in thread
From: Rainer Orth @ 2020-10-22 10:44 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: Tobias Burnus, Kwok Cheung Yeung, GCC Patches
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
--
-----------------------------------------------------------------------------
Rainer Orth, Center for Biotechnology, Bielefeld University
^ permalink raw reply [flat|nested] 13+ messages in thread
end of thread, other threads:[~2020-10-22 10:44 UTC | newest]
Thread overview: 13+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-10-14 20:32 [PATCH] openmp: Implement support for OMP_TARGET_OFFLOAD Kwok Cheung Yeung
2020-10-15 10:50 ` Jakub Jelinek
2020-10-15 11:02 ` Jakub Jelinek
2020-10-19 17:57 ` Kwok Cheung Yeung
2020-10-19 18:21 ` Jakub Jelinek
2020-10-20 12:11 ` Tobias Burnus
2020-10-20 12:17 ` Tobias Burnus
2020-10-20 12:57 ` Jakub Jelinek
2020-10-20 21:28 ` Kwok Cheung Yeung
2020-10-20 16:39 ` Rainer Orth
2020-10-20 16:45 ` Jakub Jelinek
2020-10-22 10:44 ` Rainer Orth
2020-10-21 11:52 ` [PATCH] openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirements Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).