Hi!
(This whole idea/patch still needs an overall re-work, as discussed, but
here is a small incremental improvement/bug fix.)
On Thu, 20 Aug 2015 22:52:58 +0000, Joseph Myers <jos...@codesourcery.com>
wrote:
> On Tue, 18 Aug 2015, Thomas Schwinge wrote:
> > [...] here is my current messy WIP patch [...]
> +/* List of offload targets, separated by colon. Defaults to the list
> + determined when configuring libgomp. */
> +static const char *gomp_offload_targets = OFFLOAD_TARGETS;
> +static bool gomp_offload_targets_init = false;
> +
> +/* Override the list of offload targets. This must be called early, and only
> + once. */
> +
> +void
> +GOMP_set_offload_targets (const char *offload_targets)
> +{
> + gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, offload_targets);
> +
> + /* Make sure this gets called early. */
> + assert (gomp_is_initialized == PTHREAD_ONCE_INIT);
> + /* Make sure this only gets called once. */
> + assert (!gomp_offload_targets_init);
> + gomp_offload_targets_init = true;
> + gomp_offload_targets = offload_targets;
> +}
This will obviously fail as soon as there are shared libraries involved,
compiled for offloading, which contain additional
GOMP_set_offload_targets constructor calls. Thus pushed to
openacc-gcc-7-branch:
commit 917e247055a37f912129ed545719182de0046adb
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Sun May 20 21:31:01 2018 +0200
[PR81886] Avoid "GOMP_set_offload_targets: Assertion
`!gomp_offload_targets_init' failed"
PR libgomp/81886
* openacc.h (enum acc_device_t): Add _acc_device_intel_mic,
_acc_device_hsa.
* oacc-init.c (get_openacc_name): Handle these.
(resolve_device): Debugging output.
* target.c (resolve_device, gomp_init_device)
(gomp_offload_target_available_p): Likewise.
(GOMP_set_offload_targets): Rewrite.
* testsuite/libgomp.oacc-c++/c++.exp: Provide offload target in
"-DACC_DEVICE_TYPE_host", and "-DACC_DEVICE_TYPE_nvidia".
* testsuite/libgomp.oacc-c/c.exp: Likewise.
* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
* testsuite/libgomp.oacc-c/offload-targets-1.c: New file.
* testsuite/libgomp.oacc-c/offload-targets-2.c: Likewise.
* testsuite/libgomp.oacc-c/offload-targets-3.c: Likewise.
* testsuite/libgomp.oacc-c/offload-targets-4.c: Likewise.
* testsuite/libgomp.oacc-c/offload-targets-5.c: Likewise.
* testsuite/libgomp.oacc-c/offload-targets-6.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: Adjust.
* testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
---
libgomp/ChangeLog.openacc | 34 ++++
libgomp/oacc-init.c | 7 +
libgomp/openacc.h | 2 +
libgomp/target.c | 178 +++++++++++++++++++--
libgomp/testsuite/libgomp.oacc-c++/c++.exp | 4 +-
.../libgomp.oacc-c-c++-common/acc-on-device-2.c | 2 +-
.../libgomp.oacc-c-c++-common/acc_on_device-1.c | 4 +-
.../libgomp.oacc-c-c++-common/pr85381-2.c | 3 +-
.../libgomp.oacc-c-c++-common/pr85381-3.c | 3 +-
.../libgomp.oacc-c-c++-common/pr85381-4.c | 3 +-
.../libgomp.oacc-c-c++-common/pr85381-5.c | 3 +-
.../testsuite/libgomp.oacc-c-c++-common/pr85381.c | 3 +-
.../libgomp.oacc-c-c++-common/pr85486-2.c | 3 +-
.../libgomp.oacc-c-c++-common/pr85486-3.c | 3 +-
.../testsuite/libgomp.oacc-c-c++-common/pr85486.c | 3 +-
libgomp/testsuite/libgomp.oacc-c/c.exp | 4 +-
.../testsuite/libgomp.oacc-c/offload-targets-1.c | 119 ++++++++++++++
.../testsuite/libgomp.oacc-c/offload-targets-2.c | 2 +
.../testsuite/libgomp.oacc-c/offload-targets-3.c | 10 ++
.../testsuite/libgomp.oacc-c/offload-targets-4.c | 11 ++
.../testsuite/libgomp.oacc-c/offload-targets-5.c | 10 ++
.../testsuite/libgomp.oacc-c/offload-targets-6.c | 11 ++
.../libgomp.oacc-fortran/acc_on_device-1-1.f90 | 4 +-
.../libgomp.oacc-fortran/acc_on_device-1-2.f | 4 +-
.../libgomp.oacc-fortran/acc_on_device-1-3.f | 4 +-
libgomp/testsuite/libgomp.oacc-fortran/fortran.exp | 4 +-
26 files changed, 400 insertions(+), 38 deletions(-)
diff --git libgomp/ChangeLog.openacc libgomp/ChangeLog.openacc
index d43b259..48b1f96 100644
--- libgomp/ChangeLog.openacc
+++ libgomp/ChangeLog.openacc
@@ -1,3 +1,37 @@
+2018-05-20 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR libgomp/81886
+ * openacc.h (enum acc_device_t): Add _acc_device_intel_mic,
+ _acc_device_hsa.
+ * oacc-init.c (get_openacc_name): Handle these.
+ (resolve_device): Debugging output.
+ * target.c (resolve_device, gomp_init_device)
+ (gomp_offload_target_available_p): Likewise.
+ (GOMP_set_offload_targets): Rewrite.
+ * testsuite/libgomp.oacc-c++/c++.exp: Provide offload target in
+ "-DACC_DEVICE_TYPE_host", and "-DACC_DEVICE_TYPE_nvidia".
+ * testsuite/libgomp.oacc-c/c.exp: Likewise.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+ * testsuite/libgomp.oacc-c/offload-targets-1.c: New file.
+ * testsuite/libgomp.oacc-c/offload-targets-2.c: Likewise.
+ * testsuite/libgomp.oacc-c/offload-targets-3.c: Likewise.
+ * testsuite/libgomp.oacc-c/offload-targets-4.c: Likewise.
+ * testsuite/libgomp.oacc-c/offload-targets-5.c: Likewise.
+ * testsuite/libgomp.oacc-c/offload-targets-6.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: Adjust.
+ * testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
+ * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
+
2018-05-18 Cesar Philippidis <ce...@codesourcery.com>
Backport from mainline
diff --git libgomp/oacc-init.c libgomp/oacc-init.c
index d8348c0..19c2687 100644
--- libgomp/oacc-init.c
+++ libgomp/oacc-init.c
@@ -92,6 +92,8 @@ goacc_register (struct gomp_device_descr *disp)
static const char *
get_openacc_name (const char *name)
{
+ /* not supported: _acc_device_intel_mic */
+ /* not supported: _acc_device_hsa */
if (strcmp (name, "nvptx") == 0)
return "nvidia";
else
@@ -108,6 +110,8 @@ name_of_acc_device_t (enum acc_device_t type)
case acc_device_host: return "host";
case acc_device_not_host: return "not_host";
case acc_device_nvidia: return "nvidia";
+ case /* not supported */ _acc_device_intel_mic:
+ case /* not supported */ _acc_device_hsa:
default: gomp_fatal ("unknown device type %u", (unsigned) type);
}
}
@@ -119,6 +123,8 @@ name_of_acc_device_t (enum acc_device_t type)
static struct gomp_device_descr *
resolve_device (acc_device_t d, bool fail_is_error)
{
+ gomp_debug (0, "%s (%d)\n", __FUNCTION__, (int) d);
+
acc_device_t d_arg = d;
switch (d)
@@ -203,6 +209,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
gomp_fatal ("device type %s not supported", name_of_acc_device_t (d));
}
+ gomp_debug (0, " %s: %d: %p\n", __FUNCTION__, (int) d, dispatchers[d]);
return dispatchers[d];
}
diff --git libgomp/openacc.h libgomp/openacc.h
index 102723a..3d6d57e 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -55,6 +55,8 @@ typedef enum acc_device_t {
/* acc_device_host_nonshm = 3 removed. */
acc_device_not_host = 4,
acc_device_nvidia = 5,
+ /* not supported */ _acc_device_intel_mic = 6,
+ /* not supported */ _acc_device_hsa = 7,
_ACC_device_hwm,
/* Ensure enumeration is layout compatible with int. */
_ACC_highest = __INT_MAX__,
diff --git libgomp/target.c libgomp/target.c
index aa27dc8..b5f86c8 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -108,6 +108,8 @@ gomp_get_num_devices (void)
static struct gomp_device_descr *
resolve_device (int device)
{
+ gomp_debug (0, "%s (%d)\n", __FUNCTION__, device);
+
int device_id;
if (device == GOMP_DEVICE_ICV)
{
@@ -137,6 +139,7 @@ resolve_device (int device)
&& !gomp_offload_target_available_p (devices[device_id].type))
return NULL;
+ gomp_debug (0, " %s (%d): %d\n", __FUNCTION__, device, device_id);
return &devices[device_id];
}
@@ -1883,6 +1886,9 @@ GOMP_offload_unregister (const void *host_table, int
target_type,
attribute_hidden void
gomp_init_device (struct gomp_device_descr *devicep)
{
+ gomp_debug (0, "%s (%s; %d; %d)\n", __FUNCTION__,
+ devicep->name, (int) devicep->type, devicep->target_id);
+
int i;
if (!devicep->init_device_func (devicep->target_id))
{
@@ -1946,6 +1952,8 @@ gomp_unload_device (struct gomp_device_descr *devicep)
attribute_hidden bool
gomp_offload_target_available_p (int type)
{
+ gomp_debug (0, "%s (%d)\n", __FUNCTION__, type);
+
bool available = false;
/* Has the offload target already been initialized? */
@@ -1987,6 +1995,7 @@ gomp_offload_target_available_p (int type)
gomp_mutex_unlock (®ister_lock);
}
+ gomp_debug (0, " %s (%d): %d\n", __FUNCTION__, type, (int) available);
return available;
}
@@ -3157,25 +3166,170 @@ offload_target_to_plugin_name (const char
*offload_target)
gomp_fatal ("Unknown offload target: %s", offload_target);
}
-/* List of offload targets, separated by colon. Defaults to the list
+/* List of requested offload targets, separated by colon. Defaults to the list
determined when configuring libgomp. */
static const char *gomp_offload_targets = OFFLOAD_TARGETS;
-static bool gomp_offload_targets_init = false;
+static bool gomp_offload_targets_set = false;
+static bool gomp_offload_targets_malloced = false;
-/* Override the list of offload targets. This must be called early, and only
- once. */
+/* This function frees gomp_offload_targets. */
+
+static void
+free_gomp_offload_targets (void)
+{
+ free ((char *) gomp_offload_targets);
+}
+
+/* Override the list of requested offload targets. This must be called
+ early, before gomp_target_init. */
void
GOMP_set_offload_targets (const char *offload_targets)
{
- gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, offload_targets);
-
- /* Make sure this gets called early. */
- assert (gomp_is_initialized == PTHREAD_ONCE_INIT);
- /* Make sure this only gets called once. */
- assert (!gomp_offload_targets_init);
- gomp_offload_targets_init = true;
- gomp_offload_targets = offload_targets;
+ gomp_debug (0, "%s (\"%s\"): %s\n", __FUNCTION__,
+ offload_targets, gomp_offload_targets);
+
+ /* TODO: multithreading, locking. */
+ /* TODO: this should not (sometimes) keep a copy of the offload_target
+ pointer, so that the caller knows what to expect. */
+ /* TODO: What actually is supposed to happen if some parts of a program are
+ compiled with, for example, "-foffload=disable" (that is, when called with
+ the empty string for offload_targets), and others for other actual
+ (possibly different) offload targets? */
+ if (gomp_is_initialized == PTHREAD_ONCE_INIT)
+ {
+ /* If we have not yet initialized, we capture all the offload targets
+ requested. We do not worry that the set of requested offload targets
+ vs. the set of available offload data will eventually match; any such
+ inconsistencies would be user error. (See also
+ gomp_offload_target_available_p.) */
+ if (!gomp_offload_targets_set)
+ gomp_offload_targets = offload_targets;
+ else if (gomp_offload_targets == offload_targets
+ || strcmp (gomp_offload_targets, offload_targets) == 0)
+ /* Nothing to do if the same. */;
+ else
+ {
+ /* Merge offload_targets into gomp_offload_targets. */
+ /* TODO: this could be simpler if we had the data available in a
+ different form. */
+ size_t gomp_offload_targets_len = strlen (gomp_offload_targets);
+ /* Maximum length. */
+ size_t len = (gomp_offload_targets_len + /* ":" */ 1
+ + strlen (offload_targets) + /* '\0' */ 1);
+ char *gomp_offload_targets_new = gomp_malloc (len);
+ memcpy (gomp_offload_targets_new,
+ gomp_offload_targets, gomp_offload_targets_len);
+ char *gomp_offload_targets_new_next
+ = gomp_offload_targets_new + gomp_offload_targets_len;
+ *gomp_offload_targets_new_next = '\0';
+ const char *cur = offload_targets;
+ while (*cur)
+ {
+ const char *cur_end = strchr (cur, ':');
+ /* If no other offload target following... */
+ if (cur_end == NULL)
+ /* ..., point to the terminating NUL character. */
+ cur_end = cur + strlen (cur);
+ size_t cur_len = cur_end - cur;
+
+ /* Do we already have this one listed? */
+ const char *haystack = gomp_offload_targets_new;
+ while (haystack != NULL)
+ {
+ if (strncmp (haystack, cur, cur_len) == 0)
+ break;
+ else
+ {
+ haystack = strchr (haystack, ':');
+ if (haystack != NULL)
+ haystack += /* ':' */ 1;
+ }
+ }
+ if (haystack == NULL)
+ {
+ /* Not yet listed; add it. */
+ if (gomp_offload_targets_new_next != gomp_offload_targets_new)
+ *gomp_offload_targets_new_next++ = ':';
+ assert (gomp_offload_targets_new_next + cur_len + /* '\0' */ 1
+ <= gomp_offload_targets_new + len);
+ memcpy (gomp_offload_targets_new_next, cur, cur_len);
+ gomp_offload_targets_new_next += cur_len;
+ *gomp_offload_targets_new_next = '\0';
+ }
+
+ if (*cur_end == '\0')
+ break;
+ cur = cur_end + /* : */ 1;
+ }
+
+ if (gomp_offload_targets_malloced)
+ free ((char *) gomp_offload_targets);
+ else
+ {
+ if (atexit (free_gomp_offload_targets) != 0)
+ gomp_fatal ("atexit failed");
+ }
+
+ gomp_offload_targets = gomp_offload_targets_new;
+ gomp_offload_targets_malloced = true;
+ }
+ }
+ else
+ {
+ /* If we have already initialized (which can happen only if a shared
+ library with another GOMP_set_offload_targets constructor call gets
+ loaded dynamically), and the user is now requesting offload targets
+ that were not requested previously, then we're out of luck: we can't
+ load new plugins now. Otherwise, we're all set. */
+ if (gomp_offload_targets == offload_targets
+ || strcmp (gomp_offload_targets, offload_targets) == 0)
+ /* All fine if the same. */;
+ else
+ {
+ /* Check offload_targets against gomp_offload_targets. */
+ /* TODO: this could be simpler if we had the data available in a
+ different form. */
+ const char *cur = offload_targets;
+ while (*cur)
+ {
+ const char *cur_end = strchr (cur, ':');
+ /* If no other offload target following... */
+ if (cur_end == NULL)
+ /* ..., point to the terminating NUL character. */
+ cur_end = cur + strlen (cur);
+ size_t cur_len = cur_end - cur;
+
+ /* Do we have this one listed? */
+ const char *haystack = gomp_offload_targets;
+ while (haystack != NULL)
+ {
+ if (strncmp (haystack, cur, cur_len) == 0)
+ break;
+ else
+ {
+ haystack = strchr (haystack, ':');
+ if (haystack != NULL)
+ haystack += /* ':' */ 1;
+ }
+ }
+ if (haystack == NULL)
+ {
+ /* Not listed. */
+ gomp_fatal ("Can't satisfy request for offload targets: %s;
have loaded: %s",
+ offload_targets, gomp_offload_targets);
+ }
+
+ if (*cur_end == '\0')
+ break;
+ cur = cur_end + /* : */ 1;
+ }
+ }
+ }
+ gomp_offload_targets_set = true;
+
+ gomp_debug (0, " %s (\"%s\"): %s\n", __FUNCTION__,
+ offload_targets, gomp_offload_targets);
}
/* This function initializes the runtime needed for offloading.
diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp
libgomp/testsuite/libgomp.oacc-c++/c++.exp
index 695b96d..2e17504 100644
--- libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -86,7 +86,7 @@ if { $lang_test_file_found } {
switch -glob $offload_target_openacc {
disable {
set acc_mem_shared 1
- set tagopt "-DACC_DEVICE_TYPE_host=1"
+ set tagopt "-DACC_DEVICE_TYPE_host=\"\""
}
nvptx* {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
@@ -102,7 +102,7 @@ if { $lang_test_file_found } {
lappend ALWAYS_CFLAGS
"additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common"
set acc_mem_shared 0
- set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
+ set tagopt
"-DACC_DEVICE_TYPE_nvidia=\"$offload_target_openacc\""
}
default {
set acc_mem_shared 0
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
index bfcb67d..758b1fc 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
@@ -14,7 +14,7 @@ int main ()
int expect = 1;
-#if ACC_DEVICE_TYPE_host
+#ifdef ACC_DEVICE_TYPE_host
expect = 0;
#endif
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
index 8112745..0270d06 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
@@ -37,7 +37,7 @@ main (int argc, char *argv[])
}
-#if !ACC_DEVICE_TYPE_host
+#ifndef ACC_DEVICE_TYPE_host
/* Offloaded. */
@@ -49,7 +49,7 @@ main (int argc, char *argv[])
abort ();
if (!acc_on_device (acc_device_not_host))
abort ();
-#if ACC_DEVICE_TYPE_nvidia
+#ifdef ACC_DEVICE_TYPE_nvidia
if (!acc_on_device (acc_device_nvidia))
abort ();
#else
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
index e5d02cf..6570c64 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
@@ -1,5 +1,6 @@
/* { dg-additional-options "-save-temps" } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } }
+ { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
int
main (void)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
index 7d9ba1b..c5d1c5a 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
@@ -1,5 +1,6 @@
/* { dg-additional-options "-save-temps -w" } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } }
+ { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
int a;
#pragma acc declare create(a)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
index 477297d..d955d79 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
@@ -1,5 +1,6 @@
/* { dg-additional-options "-save-temps -w" } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } }
+ { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
#define n 1024
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c
index 4653009..61e7e48 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c
@@ -1,5 +1,6 @@
/* { dg-additional-options "-save-temps" } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } }
+ { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
#define n 1024
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c
index f585ae5..2864dfc 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c
@@ -1,5 +1,6 @@
/* { dg-additional-options "-save-temps" } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } }
+ { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
int
main (void)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
index a92b5dd..0f74921 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
@@ -1,5 +1,4 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-fopenacc-dim=-:-:128" } */
/* Minimized from ref-1.C. */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
index ae62206..b4ef878 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
@@ -1,5 +1,4 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-fopenacc-dim=-:-:-" } */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "-:-:128" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
index f91dee0..99c0805 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
@@ -1,5 +1,4 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1" } } */
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* Minimized from ref-1.C. */
diff --git libgomp/testsuite/libgomp.oacc-c/c.exp
libgomp/testsuite/libgomp.oacc-c/c.exp
index 16f8295..73a7a5a 100644
--- libgomp/testsuite/libgomp.oacc-c/c.exp
+++ libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -47,7 +47,7 @@ foreach offload_target_openacc $offload_targets_s_openacc {
switch -glob $offload_target_openacc {
disable {
set acc_mem_shared 1
- set tagopt "-DACC_DEVICE_TYPE_host=1"
+ set tagopt "-DACC_DEVICE_TYPE_host=\"\""
}
nvptx* {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
@@ -63,7 +63,7 @@ foreach offload_target_openacc $offload_targets_s_openacc {
lappend ALWAYS_CFLAGS
"additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common"
set acc_mem_shared 0
- set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
+ set tagopt "-DACC_DEVICE_TYPE_nvidia=\"$offload_target_openacc\""
}
default {
set acc_mem_shared 0
diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c
libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c
new file mode 100644
index 0000000..b62a587
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c
@@ -0,0 +1,119 @@
+/* Test what happens for repeated GOMP_set_offload_targets calls, which happens
+ when shared libraries are involved, for example. As in the libgomp
+ testsuite infrastructure, it is difficult to build and link against shared
+ libraries, we simulate that by replicating some relevant
+ GOMP_set_offload_targets calls. */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <openacc.h>
+#include "libgomp_g.h"
+
+int main ()
+{
+ /* Before getting here, GOMP_set_offload_targets already got called via a
+ constructor. */
+
+ bool acc_device_types_requested[_ACC_device_hwm];
+ for (int i = 0; i < _ACC_device_hwm; ++i)
+ acc_device_types_requested[i] = false;
+
+ /* We're building for only one offload target ("-foffload=[...]") which is
+ the following. */
+ const char *offload_target_requested;
+ acc_device_t acc_device_type_requested;
+#if defined ACC_DEVICE_TYPE_nvidia
+ offload_target_requested = ACC_DEVICE_TYPE_nvidia;
+ acc_device_type_requested = acc_device_nvidia;
+#elif defined ACC_DEVICE_TYPE_host
+ offload_target_requested = ACC_DEVICE_TYPE_host;
+ acc_device_type_requested = acc_device_host;
+#else
+# error Not ported to this ACC_DEVICE_TYPE
+#endif
+ acc_device_types_requested[acc_device_type_requested] = true;
+
+#ifdef OFFLOAD_TARGETS_SAME_AGAIN
+ /* Call again; will have no noticeable difference. */
+ GOMP_set_offload_targets (offload_target_requested);
+#endif
+
+#ifdef OFFLOAD_TARGETS_ADD_EARLY
+ /* Request a (non-existing) offloading target (which will result in a
+ non-fatal diagnostic). */
+ GOMP_set_offload_targets (OFFLOAD_TARGETS_ADD);
+#endif
+
+#ifdef OFFLOAD_TARGETS_SAME_AGAIN
+ /* Call again; will have no noticeable difference. */
+ GOMP_set_offload_targets (offload_target_requested);
+ char *s;
+ {
+ size_t len = 3 * (strlen (offload_target_requested) + 1);
+# ifdef OFFLOAD_TARGETS_ADD_EARLY
+ len += 3 * (strlen (OFFLOAD_TARGETS_ADD) + 1);
+# endif
+ s = malloc (len);
+ if (s == NULL)
+ __builtin_abort ();
+ size_t len_;
+# ifndef OFFLOAD_TARGETS_ADD_EARLY
+ len_ = sprintf (s, "%s:%s:%s",
+ offload_target_requested,
+ offload_target_requested,
+ offload_target_requested);
+# else
+ len_ = sprintf (s, "%s:%s:%s:%s:%s:%s",
+ offload_target_requested,
+ offload_target_requested,
+ OFFLOAD_TARGETS_ADD,
+ OFFLOAD_TARGETS_ADD,
+ offload_target_requested,
+ OFFLOAD_TARGETS_ADD);
+# endif
+ if (len_ + 1 != len)
+ __builtin_abort ();
+ GOMP_set_offload_targets (s);
+ }
+#endif
+
+ /* Calling acc_get_num_devices will implicitly initialize offloading. */
+#if defined OFFLOAD_TARGETS_ADD_EARLY
+ fprintf (stderr, "CheCKpOInT1\n");
+#endif
+ /* acc_device_host is always available. */
+ if ((acc_get_num_devices (acc_device_host) > 0) == false)
+ __builtin_abort ();
+#if defined OFFLOAD_TARGETS_ADD_EARLY
+ fprintf (stderr, "WrONg WAy1\n");
+#endif
+ for (acc_device_t acc_device_type = acc_device_not_host + 1;
+ acc_device_type < _ACC_device_hwm;
+ ++acc_device_type)
+ {
+ /* The requested device type must be available. Any other device types
+ must not be available. */
+ if ((acc_get_num_devices (acc_device_type) > 0)
+ != acc_device_types_requested[acc_device_type])
+ __builtin_abort ();
+ }
+
+#ifdef OFFLOAD_TARGETS_SAME_AGAIN
+ /* Request the same again; will have no noticeable difference. */
+ GOMP_set_offload_targets (offload_target_requested);
+#endif
+#if defined OFFLOAD_TARGETS_ADD_LATE
+ fprintf (stderr, "CheCKpOInT2\n");
+ GOMP_set_offload_targets (OFFLOAD_TARGETS_ADD);
+ fprintf (stderr, "WrONg WAy2\n");
+#endif
+#ifdef OFFLOAD_TARGETS_SAME_AGAIN
+ GOMP_set_offload_targets (s);
+
+ /* Implementation defail: OK to "free (s)", in this case. */
+ free (s);
+#endif
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c
libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c
new file mode 100644
index 0000000..977c559
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c
@@ -0,0 +1,2 @@
+#define OFFLOAD_TARGETS_SAME_AGAIN
+#include "offload-targets-1.c"
diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c
libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c
new file mode 100644
index 0000000..1eb080b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c
@@ -0,0 +1,10 @@
+#define OFFLOAD_TARGETS_ADD "XYZ"
+#define OFFLOAD_TARGETS_ADD_EARLY
+#include "offload-targets-1.c"
+
+/*
+ { dg-output "CheCKpOInT1(\n|\r\n|\r)+" }
+ { dg-output "libgomp: Unknown offload target: XYZ(\n|\r\n|\r)+" }
+ { dg-output "$" }
+ { dg-shouldfail "" }
+*/
diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c
libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c
new file mode 100644
index 0000000..2bb7204
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c
@@ -0,0 +1,11 @@
+#define OFFLOAD_TARGETS_SAME_AGAIN
+#define OFFLOAD_TARGETS_ADD "XYZ"
+#define OFFLOAD_TARGETS_ADD_EARLY
+#include "offload-targets-1.c"
+
+/*
+ { dg-output "CheCKpOInT1(\n|\r\n|\r)+" }
+ { dg-output "libgomp: Unknown offload target: XYZ(\n|\r\n|\r)+" }
+ { dg-output "$" }
+ { dg-shouldfail "" }
+*/
diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c
libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c
new file mode 100644
index 0000000..8ba0792
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c
@@ -0,0 +1,10 @@
+#define OFFLOAD_TARGETS_ADD "XYZ"
+#define OFFLOAD_TARGETS_ADD_LATE
+#include "offload-targets-1.c"
+
+/*
+ { dg-output "CheCKpOInT2(\n|\r\n|\r)+" }
+ { dg-output "libgomp: Can't satisfy request for offload targets: XYZ; have
loaded: \[a-z-\]*(\n|\r\n|\r)+" }
+ { dg-output "$" }
+ { dg-shouldfail "" }
+*/
diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c
libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c
new file mode 100644
index 0000000..4b15582
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c
@@ -0,0 +1,11 @@
+#define OFFLOAD_TARGETS_SAME_AGAIN
+#define OFFLOAD_TARGETS_ADD "XYZ"
+#define OFFLOAD_TARGETS_ADD_LATE
+#include "offload-targets-1.c"
+
+/*
+ { dg-output "CheCKpOInT2(\n|\r\n|\r)+" }
+ { dg-output "libgomp: Can't satisfy request for offload targets: XYZ; have
loaded: \[a-z-\]*(\n|\r\n|\r)+" }
+ { dg-output "$" }
+ { dg-shouldfail "" }
+*/
diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
index 1a10f32..f57a2f2 100644
--- libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
@@ -25,7 +25,7 @@ if (acc_on_device (acc_device_nvidia)) call abort
!$acc end parallel
-#if !ACC_DEVICE_TYPE_host
+#ifndef ACC_DEVICE_TYPE_host
! Offloaded.
@@ -33,7 +33,7 @@ if (acc_on_device (acc_device_nvidia)) call abort
if (acc_on_device (acc_device_none)) call abort
if (acc_on_device (acc_device_host)) call abort
if (.not. acc_on_device (acc_device_not_host)) call abort
-#if ACC_DEVICE_TYPE_nvidia
+#ifdef ACC_DEVICE_TYPE_nvidia
if (.not. acc_on_device (acc_device_nvidia)) call abort
#else
if (acc_on_device (acc_device_nvidia)) call abort
diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
index cbd1dd9..6209d12 100644
--- libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
+++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
@@ -26,7 +26,7 @@
!$ACC END PARALLEL
-#if !ACC_DEVICE_TYPE_host
+#ifndef ACC_DEVICE_TYPE_host
! Offloaded.
@@ -34,7 +34,7 @@
IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
-#if ACC_DEVICE_TYPE_nvidia
+#ifdef ACC_DEVICE_TYPE_nvidia
IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
#else
IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
index c391776..90d567f 100644
--- libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
+++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
@@ -25,7 +25,7 @@
!$ACC END PARALLEL
-#if !ACC_DEVICE_TYPE_host
+#ifndef ACC_DEVICE_TYPE_host
! Offloaded.
@@ -33,7 +33,7 @@
IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
-#if ACC_DEVICE_TYPE_nvidia
+#ifdef ACC_DEVICE_TYPE_nvidia
IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
#else
IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
diff --git libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index d78ce55..865c704 100644
--- libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -71,7 +71,7 @@ if { $lang_test_file_found } {
switch -glob $offload_target_openacc {
disable {
set acc_mem_shared 1
- set tagopt "-DACC_DEVICE_TYPE_host=1"
+ set tagopt "-DACC_DEVICE_TYPE_host=\"\""
}
nvptx* {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
@@ -81,7 +81,7 @@ if { $lang_test_file_found } {
}
set acc_mem_shared 0
- set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
+ set tagopt
"-DACC_DEVICE_TYPE_nvidia=\"$offload_target_openacc\""
}
default {
set acc_mem_shared 0
Grüße
Thomas
