@@ -1,3 +1,37 @@
+2018-05-20 Thomas Schwinge <thomas@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 <cesar@codesourcery.com>
Backport from mainline
@@ -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];
}
@@ -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__,
@@ -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.
@@ -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
@@ -14,7 +14,7 @@ int main ()
int expect = 1;
-#if ACC_DEVICE_TYPE_host
+#ifdef ACC_DEVICE_TYPE_host
expect = 0;
#endif
@@ -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
@@ -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)
@@ -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)
@@ -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
@@ -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
@@ -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)
@@ -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. */
@@ -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" } */
@@ -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. */
@@ -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
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -0,0 +1,2 @@
+#define OFFLOAD_TARGETS_SAME_AGAIN
+#include "offload-targets-1.c"
new file mode 100644
@@ -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 "" }
+*/
new file mode 100644
@@ -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 "" }
+*/
new file mode 100644
@@ -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 "" }
+*/
new file mode 100644
@@ -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 "" }
+*/
@@ -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
@@ -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
@@ -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
@@ -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