Hi! On Tue, 29 Sep 2015 10:18:14 +0200, Jakub Jelinek <ja...@redhat.com> wrote: > On Mon, Sep 28, 2015 at 11:39:10AM +0200, Thomas Schwinge wrote: > > On Fri, 11 Sep 2015 17:43:49 +0200, Jakub Jelinek <ja...@redhat.com> wrote: > > > So, do I understand well that you'll call GOMP_set_offload_targets from > > > construct[ors] of all shared libraries (and the binary) that contain > > > offloaded > > > code? If yes, that is surely going to fail the assertions in there. > > > > Indeed. My original plan has been to generate/invoke this constructor > > only for/from the final executable and not for any shared libraries, but > > it seems I didn't implemented this correctly. > > How would you mean to implement it?
I have come to realize that we need to generate/invoke this constructor
>From everything that links against libgomp (which is what I implemented),
that is, executables as well as shared libraries.
> -fopenmp or -fopenacc code with
> offloading bits might not be in the final executable at all, nor in shared
> libraries it is linked against; such libraries could be only dlopened,
> consider say python plugin. And this is not just made up, perhaps not with
> offloading yet, but people regularly use OpenMP code in plugins and then we
> get complains that fork child of the main program is not allowed to do
> anything but async-signal-safe functions.
I'm not sure I'm completely understanding that paragraph? Are you saying
that offloaded code can be in libraries that are not linked against
libgomp? How would these register (GOMP_offload_register) their
offloaded code? I think it's a reasonable to expect that every shared
library that contains offloaded code must link against libgomp, which
will happen automatically given that it is built with -fopenmp/-fopenacc?
> > > You can dlopen such libraries etc. What if you link one library with
> > > -fopenmp=nvptx-none and another one with
> > > -fopenmp=x86_64-intelmicemul-linux?
> >
> > So, the first question to answer is: what do we expect to happen in this
> > case, or similarly, if the executable and any shared libraries are
> > compiled with different/incompatible -foffload options?
>
> As the device numbers are per-process, the only possibility I see is that
> all the physically available devices are always available, and just if you
> try to offload from some code to a device that doesn't support it, you get
> host fallback. Because, one shared library could carefully use device(xyz)
> to offload to say XeonPhi it is compiled for and supports, and another
> library device(abc) to offload to PTX it is compiled for and supports.
OK, I think I get that, and it makes sense. Even though, I don't know
how you'd do that today: as far as I can tell, there is no specification
covering the OpenMP 4 target device IDs, so I have no idea how a user
program/library could realiably use them in practice? For example, in
the current GCC implementation, the OpenMP 4 target device IDs depend on
the number of individual devices availble in the system, and the order in
which libgomp loads the plugins, which is defined (arbitrarily) by the
GCC configuration?
> > For this, I propose that the only mode of operation that we currently can
> > support is that all of the executable and any shared libraries agree on
> > the offload targets specified by -foffload, and I thus propose the
> > following patch on top of what Joseph has posted before (passes the
> > testsuite, but not yet tested otherwise):
>
> See above, no.
OK.
How's the following (complete patch instead of incremental patch; the
driver changes are still the same as before)? The changes are:
* libgomp/target.c:gomp_target_init again loads all the plugins.
* libgomp/target.c:resolve_device and
libgomp/oacc-init.c:resolve_device verify that a default device
(OpenMP device-var ICV, and acc_device_default, respectively) is
actually enabled, or resort to host fallback if not.
* GOMP_set_offload_targets renamed to GOMP_enable_offload_targets; used
to enable devices specified by -foffload. Can be called multiple
times (executable, any shared libraries); the set of enabled devices
is the union of all those ever requested.
* GOMP_offload_register (but not the new GOMP_offload_register_ver)
changed to enable all devices. This is to maintain compatibility
with old executables and shared libraries built without the -foffload
constructor support.
* IntelMIC mkoffload changed to use GOMP_offload_register_ver instead
of GOMP_offload_register, and GOMP_offload_unregister_ver instead of
GOMP_offload_unregister. To avoid enabling all devices
(GOMP_offload_register).
* New test cases to verify this (-foffload=disable, host fallback).
Ilya, I'm aware of your work on additional changes (shared memory),
<http://news.gmane.org/find-root.php?message_id=%3CCADG%3DZ0EBuhj89WEZdmaNUPy%3DE%3D63BmWofS8An8nY7rygTmdJ_w%40mail.gmail.com%3E>,
but I think my patch is still an improvement already?
Jakub, is this OK as an incremental step forward?
gcc/config/i386/intelmic-mkoffload.c | 20 +-
gcc/fortran/gfortranspec.c | 2 +-
gcc/gcc.c | 139 +++++++++++---
gcc/gcc.h | 2 +-
gcc/java/jvspec.c | 2 +-
libgomp/config.h.in | 2 +-
libgomp/configure | 6 +-
libgomp/libgomp-plugin.h | 3 +-
libgomp/libgomp.h | 1 +
libgomp/libgomp.map | 1 +
libgomp/libgomp_g.h | 1 +
libgomp/oacc-init.c | 18 +-
libgomp/plugin/configfrag.ac | 8 +-
libgomp/target.c | 205 +++++++++++++++++----
libgomp/testsuite/lib/libgomp.exp | 24 +--
.../libgomp.c++/target-1-foffload_disable.C | 3 +
.../libgomp.c++/target-foffload_disable.C | 3 +
.../libgomp.c/target-1-foffload_disable.c | 3 +
.../testsuite/libgomp.c/target-foffload_disable.c | 18 ++
.../libgomp.fortran/target-foffload_disable.f | 14 ++
.../libgomp.fortran/target1-foffload_disable.f90 | 3 +
libgomp/testsuite/libgomp.oacc-c++/c++.exp | 14 +-
libgomp/testsuite/libgomp.oacc-c/c.exp | 13 +-
libgomp/testsuite/libgomp.oacc-fortran/fortran.exp | 14 +-
24 files changed, 388 insertions(+), 131 deletions(-)
diff --git gcc/config/i386/intelmic-mkoffload.c
gcc/config/i386/intelmic-mkoffload.c
index 5195d91..e1cfbf9 100644
--- gcc/config/i386/intelmic-mkoffload.c
+++ gcc/config/i386/intelmic-mkoffload.c
@@ -360,26 +360,34 @@ generate_host_descr_file (const char *host_compiler)
"#ifdef __cplusplus\n"
"extern \"C\"\n"
"#endif\n"
- "void GOMP_offload_register (const void *, int, const void *);\n"
+ "void GOMP_offload_register_ver "
+ "(unsigned version, const void *, int, const void *);\n"
"#ifdef __cplusplus\n"
"extern \"C\"\n"
"#endif\n"
- "void GOMP_offload_unregister (const void *, int, const void *);\n\n"
+ "void GOMP_offload_unregister_ver "
+ "(unsigned version, const void *, int, const void *);\n\n"
"__attribute__((constructor))\n"
"static void\n"
"init (void)\n"
"{\n"
- " GOMP_offload_register (&__OFFLOAD_TABLE__, %d,
__offload_target_data);\n"
- "}\n\n", GOMP_DEVICE_INTEL_MIC);
+ " GOMP_offload_register_ver (%#x, &__OFFLOAD_TABLE__, "
+ "%d, __offload_target_data);\n"
+ "}\n\n",
+ GOMP_VERSION_PACK (GOMP_VERSION, GOMP_VERSION_INTEL_MIC),
+ GOMP_DEVICE_INTEL_MIC);
fprintf (src_file,
"__attribute__((destructor))\n"
"static void\n"
"fini (void)\n"
"{\n"
- " GOMP_offload_unregister (&__OFFLOAD_TABLE__, %d,
__offload_target_data);\n"
- "}\n", GOMP_DEVICE_INTEL_MIC);
+ " GOMP_offload_unregister_ver (%#x, &__OFFLOAD_TABLE__, "
+ "%d, __offload_target_data);\n"
+ "}\n",
+ GOMP_VERSION_PACK (GOMP_VERSION, GOMP_VERSION_INTEL_MIC),
+ GOMP_DEVICE_INTEL_MIC);
fclose (src_file);
diff --git gcc/fortran/gfortranspec.c gcc/fortran/gfortranspec.c
index fe594db..e3e83ba 100644
--- gcc/fortran/gfortranspec.c
+++ gcc/fortran/gfortranspec.c
@@ -439,7 +439,7 @@ int
lang_specific_pre_link (void)
{
if (library)
- do_spec ("%:include(libgfortran.spec)");
+ do_spec ("%:include(libgfortran.spec)", 0);
return 0;
}
diff --git gcc/gcc.c gcc/gcc.c
index 55a7255..ce5cba1 100644
--- gcc/gcc.c
+++ gcc/gcc.c
@@ -401,6 +401,8 @@ static const char *compare_debug_auxbase_opt_spec_function
(int, const char **);
static const char *pass_through_libs_spec_func (int, const char **);
static const char *replace_extension_spec_func (int, const char **);
static const char *greater_than_spec_func (int, const char **);
+static const char *add_omp_infile_spec_func (int, const char **);
+
static char *convert_white_space (char *);
/* The Specs Language
@@ -1189,6 +1191,11 @@ static const char *const multilib_defaults_raw[] =
MULTILIB_DEFAULTS;
static const char *const driver_self_specs[] = {
"%{fdump-final-insns:-fdump-final-insns=.} %<fdump-final-insns",
+#ifdef ENABLE_OFFLOADING
+ /* If linking against libgomp, add a setup file. */
+ "%{fopenacc|fopenmp|%:gt(%{ftree-parallelize-loops=*} 1):" \
+ "%:add-omp-infile()}",
+#endif /* ENABLE_OFFLOADING */
DRIVER_SELF_SPECS, CONFIGURE_SPECS, GOMP_SELF_SPECS, GTM_SELF_SPECS,
CILK_SELF_SPECS
};
@@ -1616,6 +1623,7 @@ static const struct spec_function static_spec_functions[]
=
{ "pass-through-libs", pass_through_libs_spec_func },
{ "replace-extension", replace_extension_spec_func },
{ "gt", greater_than_spec_func },
+ { "add-omp-infile", add_omp_infile_spec_func },
#ifdef EXTRA_SPEC_FUNCTIONS
EXTRA_SPEC_FUNCTIONS
#endif
@@ -3212,7 +3220,8 @@ execute (void)
The `validated' field describes whether any spec has looked at this switch;
if it remains false at the end of the run, the switch must be meaningless.
The `ordering' field is used to temporarily mark switches that have to be
- kept in a specific order. */
+ kept in a specific order.
+ The `lang_mask' field stores the flags associated with this option. */
#define SWITCH_LIVE (1 << 0)
#define SWITCH_FALSE (1 << 1)
@@ -3228,6 +3237,7 @@ struct switchstr
bool known;
bool validated;
bool ordering;
+ unsigned int lang_mask;
};
static struct switchstr *switches;
@@ -3236,6 +3246,10 @@ static int n_switches;
static int n_switches_alloc;
+/* If nonzero, do not pass through switches for languages not matching
+ this mask. */
+static unsigned int spec_lang_mask_accept;
+
/* Set to zero if -fcompare-debug is disabled, positive if it's
enabled and we're running the first compilation, negative if it's
enabled and we're running the second compilation. For most of the
@@ -3273,6 +3287,7 @@ struct infile
const char *name;
const char *language;
struct compiler *incompiler;
+ unsigned int lang_mask;
bool compiled;
bool preprocessed;
};
@@ -3466,15 +3481,16 @@ alloc_infile (void)
}
}
-/* Store an input file with the given NAME and LANGUAGE in
+/* Store an input file with the given NAME and LANGUAGE and LANG_MASK in
infiles. */
static void
-add_infile (const char *name, const char *language)
+add_infile (const char *name, const char *language, unsigned int lang_mask)
{
alloc_infile ();
infiles[n_infiles].name = name;
- infiles[n_infiles++].language = language;
+ infiles[n_infiles].language = language;
+ infiles[n_infiles++].lang_mask = lang_mask;
}
/* Allocate space for a switch in switches. */
@@ -3495,11 +3511,12 @@ alloc_switch (void)
}
/* Save an option OPT with N_ARGS arguments in array ARGS, marking it
- as validated if VALIDATED and KNOWN if it is an internal switch. */
+ as validated if VALIDATED and KNOWN if it is an internal switch.
+ LANG_MASK is the flags associated with this option. */
static void
save_switch (const char *opt, size_t n_args, const char *const *args,
- bool validated, bool known)
+ bool validated, bool known, unsigned int lang_mask)
{
alloc_switch ();
switches[n_switches].part1 = opt + 1;
@@ -3516,6 +3533,7 @@ save_switch (const char *opt, size_t n_args, const char
*const *args,
switches[n_switches].validated = validated;
switches[n_switches].known = known;
switches[n_switches].ordering = 0;
+ switches[n_switches].lang_mask = lang_mask;
n_switches++;
}
@@ -3533,7 +3551,8 @@ driver_unknown_option_callback (const struct
cl_decoded_option *decoded)
diagnosed only if there are warnings. */
save_switch (decoded->canonical_option[0],
decoded->canonical_option_num_elements - 1,
- &decoded->canonical_option[1], false, true);
+ &decoded->canonical_option[1], false, true,
+ cl_options[decoded->opt_index].flags);
return false;
}
if (decoded->opt_index == OPT_SPECIAL_unknown)
@@ -3541,7 +3560,8 @@ driver_unknown_option_callback (const struct
cl_decoded_option *decoded)
/* Give it a chance to define it a spec file. */
save_switch (decoded->canonical_option[0],
decoded->canonical_option_num_elements - 1,
- &decoded->canonical_option[1], false, false);
+ &decoded->canonical_option[1], false, false,
+ cl_options[decoded->opt_index].flags);
return false;
}
else
@@ -3568,7 +3588,8 @@ driver_wrong_lang_callback (const struct
cl_decoded_option *decoded,
else
save_switch (decoded->canonical_option[0],
decoded->canonical_option_num_elements - 1,
- &decoded->canonical_option[1], false, true);
+ &decoded->canonical_option[1], false, true,
+ option->flags);
}
static const char *spec_lang = 0;
@@ -3817,7 +3838,8 @@ driver_handle_option (struct gcc_options *opts,
compare_debug_opt = NULL;
else
compare_debug_opt = arg;
- save_switch (compare_debug_replacement_opt, 0, NULL, validated, true);
+ save_switch (compare_debug_replacement_opt, 0, NULL, validated, true,
+ cl_options[opt_index].flags);
return true;
case OPT_fdiagnostics_color_:
@@ -3872,17 +3894,17 @@ driver_handle_option (struct gcc_options *opts,
for (j = 0; arg[j]; j++)
if (arg[j] == ',')
{
- add_infile (save_string (arg + prev, j - prev), "*");
+ add_infile (save_string (arg + prev, j - prev), "*", 0);
prev = j + 1;
}
/* Record the part after the last comma. */
- add_infile (arg + prev, "*");
+ add_infile (arg + prev, "*", 0);
}
do_save = false;
break;
case OPT_Xlinker:
- add_infile (arg, "*");
+ add_infile (arg, "*", 0);
do_save = false;
break;
@@ -3899,19 +3921,21 @@ driver_handle_option (struct gcc_options *opts,
case OPT_l:
/* POSIX allows separation of -l and the lib arg; canonicalize
by concatenating -l with its arg */
- add_infile (concat ("-l", arg, NULL), "*");
+ add_infile (concat ("-l", arg, NULL), "*", 0);
do_save = false;
break;
case OPT_L:
/* Similarly, canonicalize -L for linkers that may not accept
separate arguments. */
- save_switch (concat ("-L", arg, NULL), 0, NULL, validated, true);
+ save_switch (concat ("-L", arg, NULL), 0, NULL, validated, true,
+ cl_options[opt_index].flags);
return true;
case OPT_F:
/* Likewise -F. */
- save_switch (concat ("-F", arg, NULL), 0, NULL, validated, true);
+ save_switch (concat ("-F", arg, NULL), 0, NULL, validated, true,
+ cl_options[opt_index].flags);
return true;
case OPT_save_temps:
@@ -4034,7 +4058,8 @@ driver_handle_option (struct gcc_options *opts,
save_temps_prefix = xstrdup (arg);
/* On some systems, ld cannot handle "-o" without a space. So
split the option from its argument. */
- save_switch ("-o", 1, &arg, validated, true);
+ save_switch ("-o", 1, &arg, validated, true,
+ cl_options[opt_index].flags);
return true;
#ifdef ENABLE_DEFAULT_PIE
@@ -4070,7 +4095,8 @@ driver_handle_option (struct gcc_options *opts,
if (do_save)
save_switch (decoded->canonical_option[0],
decoded->canonical_option_num_elements - 1,
- &decoded->canonical_option[1], validated, true);
+ &decoded->canonical_option[1], validated, true,
+ cl_options[opt_index].flags);
return true;
}
@@ -4367,7 +4393,7 @@ process_command (unsigned int decoded_options_count,
if (strcmp (fname, "-") != 0 && access (fname, F_OK) < 0)
perror_with_name (fname);
else
- add_infile (arg, spec_lang);
+ add_infile (arg, spec_lang, 0);
free (fname);
continue;
@@ -4516,7 +4542,8 @@ process_command (unsigned int decoded_options_count,
if (compare_debug == 2 || compare_debug == 3)
{
const char *opt = concat ("-fcompare-debug=", compare_debug_opt, NULL);
- save_switch (opt, 0, NULL, false, true);
+ save_switch (opt, 0, NULL, false, true,
+ cl_options[OPT_fcompare_debug_].flags);
compare_debug = 1;
}
@@ -4527,7 +4554,7 @@ process_command (unsigned int decoded_options_count,
/* Create a dummy input file, so that we can pass
the help option on to the various sub-processes. */
- add_infile ("help-dummy", "c");
+ add_infile ("help-dummy", "c", 0);
}
alloc_switch ();
@@ -4728,13 +4755,15 @@ insert_wrapper (const char *wrapper)
}
/* Process the spec SPEC and run the commands specified therein.
+ If LANG_MASK is nonzero, switches for other languages are discarded.
Returns 0 if the spec is successfully processed; -1 if failed. */
int
-do_spec (const char *spec)
+do_spec (const char *spec, unsigned int lang_mask)
{
int value;
+ spec_lang_mask_accept = lang_mask;
value = do_spec_2 (spec);
/* Force out any unfinished command.
@@ -4892,7 +4921,8 @@ do_self_spec (const char *spec)
save_switch (decoded_options[j].canonical_option[0],
(decoded_options[j].canonical_option_num_elements
- 1),
- &decoded_options[j].canonical_option[1], false,
true);
+ &decoded_options[j].canonical_option[1], false, true,
+ cl_options[decoded_options[j].opt_index].flags);
break;
default:
@@ -6488,6 +6518,14 @@ check_live_switch (int switchnum, int prefix_length)
static void
give_switch (int switchnum, int omit_first_word)
{
+ int lang_mask = switches[switchnum].lang_mask & ((1U << cl_lang_count) - 1);
+ unsigned int lang_mask_accept = (1U << cl_lang_count) - 1;
+ if (spec_lang_mask_accept != 0)
+ lang_mask_accept = spec_lang_mask_accept;
+ /* Drop switches specific to a language not in the given mask. */
+ if (lang_mask != 0 && !(lang_mask & lang_mask_accept))
+ return;
+
if ((switches[switchnum].live_cond & SWITCH_IGNORE) != 0)
return;
@@ -7589,9 +7627,6 @@ driver::maybe_putenv_OFFLOAD_TARGETS () const
strlen (offload_targets) + 1);
xputenv (XOBFINISH (&collect_obstack, char *));
}
-
- free (offload_targets);
- offload_targets = NULL;
}
/* Reject switches that no pass was interested in. */
@@ -7895,7 +7930,8 @@ driver::do_spec_on_infiles () const
debug_check_temp_file[1] = NULL;
}
- value = do_spec (input_file_compiler->spec);
+ value = do_spec (input_file_compiler->spec,
+ infiles[i].lang_mask);
infiles[i].compiled = true;
if (value < 0)
this_file_error = 1;
@@ -7909,7 +7945,8 @@ driver::do_spec_on_infiles () const
n_switches_alloc = n_switches_alloc_debug_check[1];
switches = switches_debug_check[1];
- value = do_spec (input_file_compiler->spec);
+ value = do_spec (input_file_compiler->spec,
+ infiles[i].lang_mask);
compare_debug = -compare_debug;
n_switches = n_switches_debug_check[0];
@@ -8064,7 +8101,7 @@ driver::maybe_run_linker (const char *argv0) const
" to the linker.\n\n"));
fflush (stdout);
}
- int value = do_spec (link_command_spec);
+ int value = do_spec (link_command_spec, 0);
if (value < 0)
errorcount = 1;
linker_was_run = (tmp != execution_count);
@@ -9655,6 +9692,50 @@ greater_than_spec_func (int argc, const char **argv)
return NULL;
}
+/* If applicable, generate a C source file containing a constructor call to
+ GOMP_enable_offload_targets, to inform libgomp which offload targets have
+ actually been requested (-foffload=[...]), and add that as an infile. */
+
+static const char *
+add_omp_infile_spec_func (int argc, const char **)
+{
+ gcc_assert (argc == 0);
+ gcc_assert (offload_targets != NULL);
+
+ /* Nothing to do if we're not actually linking. */
+ if (have_c)
+ return NULL;
+
+ int err;
+ const char *tmp_filename;
+ tmp_filename = make_temp_file (".c");
+ record_temp_file (tmp_filename, !save_temps_flag, 0);
+ FILE *f = fopen (tmp_filename, "w");
+ if (f == NULL)
+ fatal_error (input_location,
+ "could not open temporary file %s", tmp_filename);
+ /* As libgomp uses constructors internally, and this code is only added when
+ linking against libgomp, it is fine to use a constructor here. */
+ err = fprintf (f,
+ "extern void GOMP_enable_offload_targets (const char *);\n"
+ "static __attribute__ ((constructor)) void\n"
+ "init (void)\n"
+ "{\n"
+ " GOMP_enable_offload_targets (\"%s\");\n"
+ "}\n",
+ offload_targets);
+ if (err < 0)
+ fatal_error (input_location,
+ "could not write to temporary file %s", tmp_filename);
+ err = fclose (f);
+ if (err == EOF)
+ fatal_error (input_location,
+ "could not close temporary file %s", tmp_filename);
+
+ add_infile (tmp_filename, "cpp-output", CL_C);
+ return NULL;
+}
+
/* Insert backslash before spaces in ORIG (usually a file path), to
avoid being broken by spec parser.
diff --git gcc/gcc.h gcc/gcc.h
index e1abe43..c71582d 100644
--- gcc/gcc.h
+++ gcc/gcc.h
@@ -68,7 +68,7 @@ struct spec_function
};
/* These are exported by gcc.c. */
-extern int do_spec (const char *);
+extern int do_spec (const char *, unsigned int);
extern void record_temp_file (const char *, int, int);
extern void pfatal_with_name (const char *) ATTRIBUTE_NORETURN;
extern void set_input (const char *);
diff --git gcc/java/jvspec.c gcc/java/jvspec.c
index d4efb73..518aa4d 100644
--- gcc/java/jvspec.c
+++ gcc/java/jvspec.c
@@ -629,7 +629,7 @@ lang_specific_pre_link (void)
class name. Append dummy `.c' that can be stripped by set_input so %b
is correct. */
set_input (concat (main_class_name, "main.c", NULL));
- err = do_spec (jvgenmain_spec);
+ err = do_spec (jvgenmain_spec, 0);
if (err == 0)
{
/* Shift the outfiles array so the generated main comes first.
diff --git libgomp/config.h.in libgomp/config.h.in
index 2e4c698..d63e56a 100644
--- libgomp/config.h.in
+++ libgomp/config.h.in
@@ -95,7 +95,7 @@
*/
#undef LT_OBJDIR
-/* Define to offload targets, separated by commas. */
+/* Define to offload targets, separated by colons. */
#undef OFFLOAD_TARGETS
/* Name of package */
diff --git libgomp/configure libgomp/configure
index 74d4e82..36ae548 100755
--- libgomp/configure
+++ libgomp/configure
@@ -15236,10 +15236,8 @@ if test x"$enable_offload_targets" != x; then
tgt=`echo $tgt | sed 's/=.*//'`
case $tgt in
*-intelmic-* | *-intelmicemul-*)
- tgt_name=intelmic
;;
nvptx*)
- tgt_name=nvptx
PLUGIN_NVPTX=$tgt
PLUGIN_NVPTX_CPPFLAGS=$CUDA_DRIVER_CPPFLAGS
PLUGIN_NVPTX_LDFLAGS=$CUDA_DRIVER_LDFLAGS
@@ -15282,9 +15280,9 @@ rm -f core conftest.err conftest.$ac_objext \
;;
esac
if test x"$offload_targets" = x; then
- offload_targets=$tgt_name
+ offload_targets=$tgt
else
- offload_targets=$offload_targets,$tgt_name
+ offload_targets=$offload_targets:$tgt
fi
if test x"$tgt_dir" != x; then
offload_additional_options="$offload_additional_options
-B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
diff --git libgomp/libgomp-plugin.h libgomp/libgomp-plugin.h
index 24fbb94..5da4fa7 100644
--- libgomp/libgomp-plugin.h
+++ libgomp/libgomp-plugin.h
@@ -48,7 +48,8 @@ enum offload_target_type
OFFLOAD_TARGET_TYPE_HOST = 2,
/* OFFLOAD_TARGET_TYPE_HOST_NONSHM = 3 removed. */
OFFLOAD_TARGET_TYPE_NVIDIA_PTX = 5,
- OFFLOAD_TARGET_TYPE_INTEL_MIC = 6
+ OFFLOAD_TARGET_TYPE_INTEL_MIC = 6,
+ OFFLOAD_TARGET_TYPE_HWM
};
/* Auxiliary struct, used for transferring pairs of addresses from plugin
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 04262c4..61cb9f3 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -636,6 +636,7 @@ extern void gomp_free_thread (void *);
extern void gomp_init_targets_once (void);
extern int gomp_get_num_devices (void);
+extern bool gomp_offload_target_enabled_p (enum offload_target_type);
typedef struct splay_tree_node_s *splay_tree_node;
typedef struct splay_tree_s *splay_tree;
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 3b3e0c2..5d65dc1 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -236,6 +236,7 @@ GOMP_4.0.1 {
GOMP_4.0.2 {
global:
+ GOMP_enable_offload_targets;
GOMP_offload_register_ver;
GOMP_offload_unregister_ver;
} GOMP_4.0.1;
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index e7f4eff..d8e7149 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -206,6 +206,7 @@ extern void GOMP_single_copy_end (void *);
/* target.c */
+extern void GOMP_enable_offload_targets (const char *);
extern void GOMP_target (int, void (*) (void *), const void *,
size_t, void **, size_t *, unsigned char *);
extern void GOMP_target_data (int, const void *,
diff --git libgomp/oacc-init.c libgomp/oacc-init.c
index 28b9e7a..c4114cc 100644
--- libgomp/oacc-init.c
+++ libgomp/oacc-init.c
@@ -122,7 +122,9 @@ resolve_device (acc_device_t d, bool fail_is_error)
{
if (goacc_device_type)
{
- /* Lookup the named device. */
+ /* Lookup the device that has been explicitly named, so do not pay
+ attention to gomp_offload_target_enabled_p. (That is, hard
+ error if not actually enabled.) */
while (++d != _ACC_device_hwm)
if (dispatchers[d]
&& !strcasecmp (goacc_device_type,
@@ -148,8 +150,14 @@ resolve_device (acc_device_t d, bool fail_is_error)
case acc_device_not_host:
/* Find the first available device after acc_device_not_host. */
while (++d != _ACC_device_hwm)
- if (dispatchers[d] && dispatchers[d]->get_num_devices_func () > 0)
+ if (dispatchers[d]
+ && dispatchers[d]->get_num_devices_func () > 0
+ /* No device has been explicitly named, so pay attention to
+ gomp_offload_target_enabled_p, to not decide on an offload
+ target that has not been enabled. */
+ && gomp_offload_target_enabled_p (dispatchers[d]->type))
goto found;
+ /* No non-host device found. */
if (d_arg == acc_device_default)
{
d = acc_device_host;
@@ -164,9 +172,6 @@ resolve_device (acc_device_t d, bool fail_is_error)
return NULL;
break;
- case acc_device_host:
- break;
-
default:
if (d > _ACC_device_hwm)
{
@@ -181,7 +186,8 @@ resolve_device (acc_device_t d, bool fail_is_error)
assert (d != acc_device_none
&& d != acc_device_default
- && d != acc_device_not_host);
+ && d != acc_device_not_host
+ && d < _ACC_device_hwm);
if (dispatchers[d] == NULL && fail_is_error)
{
diff --git libgomp/plugin/configfrag.ac libgomp/plugin/configfrag.ac
index ad70dd1..a1bfec6 100644
--- libgomp/plugin/configfrag.ac
+++ libgomp/plugin/configfrag.ac
@@ -92,10 +92,8 @@ if test x"$enable_offload_targets" != x; then
tgt=`echo $tgt | sed 's/=.*//'`
case $tgt in
*-intelmic-* | *-intelmicemul-*)
- tgt_name=intelmic
;;
nvptx*)
- tgt_name=nvptx
PLUGIN_NVPTX=$tgt
PLUGIN_NVPTX_CPPFLAGS=$CUDA_DRIVER_CPPFLAGS
PLUGIN_NVPTX_LDFLAGS=$CUDA_DRIVER_LDFLAGS
@@ -127,9 +125,9 @@ if test x"$enable_offload_targets" != x; then
;;
esac
if test x"$offload_targets" = x; then
- offload_targets=$tgt_name
+ offload_targets=$tgt
else
- offload_targets=$offload_targets,$tgt_name
+ offload_targets=$offload_targets:$tgt
fi
if test x"$tgt_dir" != x; then
offload_additional_options="$offload_additional_options
-B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
@@ -141,7 +139,7 @@ if test x"$enable_offload_targets" != x; then
done
fi
AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets",
- [Define to offload targets, separated by commas.])
+ [Define to offload targets, separated by colons.])
AM_CONDITIONAL([PLUGIN_NVPTX], [test $PLUGIN_NVPTX = 1])
AC_DEFINE_UNQUOTED([PLUGIN_NVPTX], [$PLUGIN_NVPTX],
[Define to 1 if the NVIDIA plugin is built, 0 if not.])
diff --git libgomp/target.c libgomp/target.c
index 6f0a339..1c8f337 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -71,6 +71,9 @@ static int num_offload_images;
/* Array of descriptors for all available devices. */
static struct gomp_device_descr *devices;
+/* Set of enabled devices. */
+static bool devices_enabled[OFFLOAD_TARGET_TYPE_HWM];
+
/* Total number of available devices. */
static int num_devices;
@@ -124,19 +127,30 @@ gomp_get_num_devices (void)
}
static struct gomp_device_descr *
-resolve_device (int device_id)
+resolve_device (int device)
{
- if (device_id == GOMP_DEVICE_ICV)
+ int device_id;
+ if (device == GOMP_DEVICE_ICV)
{
struct gomp_task_icv *icv = gomp_icv (false);
device_id = icv->default_device_var;
}
+ else
+ device_id = device;
if (device_id < 0 || device_id >= gomp_get_num_devices ())
return NULL;
/* As it is immutable once it has been initialized, it's safe to access
devices without register_lock held. */
+
+ /* If the device specified by the device-var ICV is not actually enabled,
+ don't try use it (which will fail if it doesn't have offload data
+ available), and use host fallback instead. */
+ if (device == GOMP_DEVICE_ICV
+ && !gomp_offload_target_enabled_p (devices[device_id].type))
+ return NULL;
+
return &devices[device_id];
}
@@ -799,6 +813,8 @@ void
GOMP_offload_register_ver (unsigned version, const void *host_table,
int target_type, const void *target_data)
{
+ gomp_debug(0, "%s (%#x, %d)\n", __FUNCTION__, version, target_type);
+
int i;
if (GOMP_VERSION_LIB (version) > GOMP_VERSION)
@@ -836,6 +852,18 @@ void
GOMP_offload_register (const void *host_table, int target_type,
const void *target_data)
{
+ gomp_debug(0, "%s (%d)\n", __FUNCTION__, target_type);
+
+ gomp_mutex_lock (®ister_lock);
+ /* If we're seeing this function called, then default to the old behavior of
+ enabling all offload targets: this is what old executables and shared
+ libraries expect. */
+ for (enum offload_target_type type = 0;
+ type < OFFLOAD_TARGET_TYPE_HWM;
+ ++type)
+ devices_enabled[type] = true;
+ gomp_mutex_unlock (®ister_lock);
+
GOMP_offload_register_ver (0, host_table, target_type, target_data);
}
@@ -847,6 +875,8 @@ void
GOMP_offload_unregister_ver (unsigned version, const void *host_table,
int target_type, const void *target_data)
{
+ gomp_debug(0, "%s (%#x, %d)\n", __FUNCTION__, version, target_type);
+
int i;
gomp_mutex_lock (®ister_lock);
@@ -877,6 +907,8 @@ void
GOMP_offload_unregister (const void *host_table, int target_type,
const void *target_data)
{
+ gomp_debug(0, "%s (%d)\n", __FUNCTION__, target_type);
+
GOMP_offload_unregister_ver (0, host_table, target_type, target_data);
}
@@ -952,6 +984,18 @@ gomp_fini_device (struct gomp_device_descr *devicep)
devicep->is_initialized = false;
}
+/* Has the offload target type TYPE been enabled?
+
+ We cannot verify that *all* offload data is available that could possibly be
+ required, so if we later find any offload data missing for this offload
+ target, then that's user error. */
+
+attribute_hidden bool
+gomp_offload_target_enabled_p (enum offload_target_type type)
+{
+ return devices_enabled[type];
+}
+
/* Called when encountering a target directive. If DEVICE
is GOMP_DEVICE_ICV, it means use device-var ICV. If it is
GOMP_DEVICE_HOST_FALLBACK (or any value
@@ -1121,6 +1165,8 @@ static bool
gomp_load_plugin_for_device (struct gomp_device_descr *device,
const char *plugin_name)
{
+ gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, plugin_name);
+
const char *err = NULL, *last_missing = NULL;
void *plugin_handle = dlopen (plugin_name, RTLD_LAZY);
@@ -1216,6 +1262,78 @@ gomp_load_plugin_for_device (struct gomp_device_descr
*device,
return 0;
}
+/* Return the corresponding offload target type for the offload target name
+ OFFLOAD_TARGET, or 0 if unknown. */
+
+static enum offload_target_type
+offload_target_to_type (const char *offload_target)
+{
+ if (strstr (offload_target, "-intelmic") != NULL)
+ return OFFLOAD_TARGET_TYPE_INTEL_MIC;
+ else if (strncmp (offload_target, "nvptx", 5) == 0)
+ return OFFLOAD_TARGET_TYPE_NVIDIA_PTX;
+ else
+ return 0;
+}
+
+/* Return the corresponding plugin name for the offload target type TYPE, or
+ NULL if unknown. */
+
+static const char *
+offload_target_type_to_plugin_name (enum offload_target_type type)
+{
+ switch (type)
+ {
+ case OFFLOAD_TARGET_TYPE_INTEL_MIC:
+ return "intelmic";
+ case OFFLOAD_TARGET_TYPE_NVIDIA_PTX:
+ return "nvptx";
+ default:
+ return NULL;
+ }
+}
+
+/* Enable the specified OFFLOAD_TARGETS, the set passed to the compiler at link
+ time. */
+
+void
+GOMP_enable_offload_targets (const char *offload_targets)
+{
+ gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, offload_targets);
+
+ char *offload_targets_dup = strdup (offload_targets);
+ if (offload_targets_dup == NULL)
+ gomp_fatal ("Out of memory");
+
+ gomp_mutex_lock (®ister_lock);
+
+ char *cur = offload_targets_dup;
+ while (cur)
+ {
+ char *next = strchr (cur, ':');
+ if (next != NULL)
+ {
+ *next = '\0';
+ ++next;
+ }
+ enum offload_target_type type = offload_target_to_type (cur);
+ if (type == 0)
+ {
+ /* An unknown offload target has been requested; ignore it. This
+ makes us (future-)proof if offload targets are requested that
+ are not supported in this build of libgomp. */
+ }
+ else
+ devices_enabled[type] = true;
+
+ cur = next;
+ }
+
+ gomp_mutex_unlock (®ister_lock);
+
+ free (offload_targets_dup);
+}
+
/* This function initializes the runtime needed for offloading.
It parses the list of offload targets and tries to load the plugins for
these targets. On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
@@ -1223,13 +1341,13 @@ gomp_load_plugin_for_device (struct gomp_device_descr
*device,
corresponding devices, first the GOMP_OFFLOAD_CAP_OPENMP_400 ones, follows
by the others. */
+static const char *gomp_plugin_prefix ="libgomp-plugin-";
+static const char *gomp_plugin_suffix = SONAME_SUFFIX (1);
+
static void
gomp_target_init (void)
{
- const char *prefix ="libgomp-plugin-";
- const char *suffix = SONAME_SUFFIX (1);
const char *cur, *next;
- char *plugin_name;
int i, new_num_devices;
gomp_mutex_lock (®ister_lock);
@@ -1241,44 +1359,58 @@ gomp_target_init (void)
if (*cur)
do
{
- struct gomp_device_descr current_device;
-
- next = strchr (cur, ',');
-
- plugin_name = (char *) malloc (1 + (next ? next - cur : strlen (cur))
- + strlen (prefix) + strlen (suffix));
- if (!plugin_name)
- {
- num_devices = 0;
- break;
- }
-
- strcpy (plugin_name, prefix);
- strncat (plugin_name, cur, next ? next - cur : strlen (cur));
- strcat (plugin_name, suffix);
+ next = strchr (cur, ':');
+ /* If no other offload target following... */
+ if (next == NULL)
+ /* ..., point to the terminating NUL character. */
+ next = strchr (cur, '\0');
+
+ size_t gomp_plugin_prefix_len = strlen (gomp_plugin_prefix);
+ size_t cur_len = next - cur;
+ size_t gomp_plugin_suffix_len = strlen (gomp_plugin_suffix);
+ char *plugin_name
+ = gomp_realloc_unlock (NULL, (gomp_plugin_prefix_len
+ + cur_len
+ + gomp_plugin_suffix_len
+ + 1));
+ memcpy (plugin_name, gomp_plugin_prefix, gomp_plugin_prefix_len);
+ memcpy (plugin_name + gomp_plugin_prefix_len, cur, cur_len);
+ /* NUL-terminate the string here... */
+ plugin_name[gomp_plugin_prefix_len + cur_len] = '\0';
+ /* ..., so that we can then use it to translate the offload target to
+ the plugin name... */
+ enum offload_target_type type
+ = offload_target_to_type (plugin_name + gomp_plugin_prefix_len);
+ const char *cur_plugin_name
+ = offload_target_type_to_plugin_name (type);
+ size_t cur_plugin_name_len = strlen (cur_plugin_name);
+ assert (cur_plugin_name_len <= cur_len);
+ /* ..., and then rewrite it. */
+ memcpy (plugin_name + gomp_plugin_prefix_len,
+ cur_plugin_name, cur_plugin_name_len);
+ memcpy (plugin_name + gomp_plugin_prefix_len + cur_plugin_name_len,
+ gomp_plugin_suffix, gomp_plugin_suffix_len);
+ plugin_name[gomp_plugin_prefix_len
+ + cur_plugin_name_len
+ + gomp_plugin_suffix_len] = '\0';
+ struct gomp_device_descr current_device;
if (gomp_load_plugin_for_device (¤t_device, plugin_name))
{
new_num_devices = current_device.get_num_devices_func ();
if (new_num_devices >= 1)
{
- /* Augment DEVICES and NUM_DEVICES. */
-
- devices = realloc (devices, (num_devices + new_num_devices)
- * sizeof (struct gomp_device_descr));
- if (!devices)
- {
- num_devices = 0;
- free (plugin_name);
- break;
- }
-
current_device.name = current_device.get_name_func ();
/* current_device.capabilities has already been set. */
current_device.type = current_device.get_type_func ();
current_device.mem_map.root = NULL;
current_device.is_initialized = false;
current_device.openacc.data_environ = NULL;
+
+ /* Augment DEVICES and NUM_DEVICES. */
+ devices = gomp_realloc_unlock
+ (devices, ((num_devices + new_num_devices)
+ * sizeof (struct gomp_device_descr)));
for (i = 0; i < new_num_devices; i++)
{
current_device.target_id = i;
@@ -1292,18 +1424,13 @@ gomp_target_init (void)
free (plugin_name);
cur = next + 1;
}
- while (next);
+ while (*next);
/* 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)
- {
- num_devices = 0;
- free (devices);
- devices = NULL;
- }
+ = gomp_realloc_unlock (NULL,
+ num_devices * sizeof (struct gomp_device_descr));
num_devices_openmp = 0;
for (i = 0; i < num_devices; i++)
if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
diff --git libgomp/testsuite/lib/libgomp.exp libgomp/testsuite/lib/libgomp.exp
index f04b163..3b4c515 100644
--- libgomp/testsuite/lib/libgomp.exp
+++ libgomp/testsuite/lib/libgomp.exp
@@ -36,24 +36,21 @@ load_gcc_lib fortran-modules.exp
load_file libgomp-test-support.exp
# Populate offload_targets_s (offloading targets separated by a space), and
-# offload_targets_s_openacc (the same, but with OpenACC names; OpenACC spells
-# some of them a little differently).
-set offload_targets_s [split $offload_targets ","]
+# offload_targets_s_openacc (those suitable for OpenACC).
+set offload_targets_s [split $offload_targets ":"]
set offload_targets_s_openacc {}
foreach offload_target_openacc $offload_targets_s {
- switch $offload_target_openacc {
- intelmic {
+ switch -glob $offload_target_openacc {
+ *-intelmic* {
# Skip; will all FAIL because of missing
# GOMP_OFFLOAD_CAP_OPENACC_200.
continue
}
- nvptx {
- set offload_target_openacc "nvidia"
- }
}
lappend offload_targets_s_openacc "$offload_target_openacc"
}
-lappend offload_targets_s_openacc "host"
+# Host fallback.
+lappend offload_targets_s_openacc "disable"
set dg-do-what-default run
@@ -134,7 +131,7 @@ proc libgomp_init { args } {
# Add liboffloadmic build directory in LD_LIBRARY_PATH to support
# non-fallback testing for Intel MIC targets
global offload_targets
- if { [string match "*,intelmic,*" ",$offload_targets,"] } {
+ if { [string match "*:*-intelmic*:*" ":$offload_targets:"] } {
append always_ld_library_path ":${blddir}/../liboffloadmic/.libs"
append always_ld_library_path ":${blddir}/../liboffloadmic/plugin/.libs"
# libstdc++ is required by liboffloadmic
@@ -332,15 +329,14 @@ proc check_effective_target_openacc_nvidia_accel_present
{ } {
}
# Return 1 if at least one nvidia board is present, and the nvidia device type
-# is selected by default by means of setting the environment variable
-# ACC_DEVICE_TYPE.
+# is selected by default.
proc check_effective_target_openacc_nvidia_accel_selected { } {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
return 0;
}
global offload_target_openacc
- if { $offload_target_openacc == "nvidia" } {
+ if { [string match "nvptx*" $offload_target_openacc] } {
return 1;
}
return 0;
@@ -350,7 +346,7 @@ proc check_effective_target_openacc_nvidia_accel_selected {
} {
proc check_effective_target_openacc_host_selected { } {
global offload_target_openacc
- if { $offload_target_openacc == "host" } {
+ if { $offload_target_openacc == "disable" } {
return 1;
}
return 0;
diff --git libgomp/testsuite/libgomp.c++/target-1-foffload_disable.C
libgomp/testsuite/libgomp.c++/target-1-foffload_disable.C
new file mode 100644
index 0000000..15b9432
--- /dev/null
+++ libgomp/testsuite/libgomp.c++/target-1-foffload_disable.C
@@ -0,0 +1,3 @@
+/* { dg-additional-options "-foffload=disable" } */
+
+#include "target-1.C"
diff --git libgomp/testsuite/libgomp.c++/target-foffload_disable.C
libgomp/testsuite/libgomp.c++/target-foffload_disable.C
new file mode 100644
index 0000000..c07dea1
--- /dev/null
+++ libgomp/testsuite/libgomp.c++/target-foffload_disable.C
@@ -0,0 +1,3 @@
+/* { dg-additional-options "-foffload=disable" } */
+
+#include "../libgomp.c/target-foffload_disable.c"
diff --git libgomp/testsuite/libgomp.c/target-1-foffload_disable.c
libgomp/testsuite/libgomp.c/target-1-foffload_disable.c
new file mode 100644
index 0000000..177cceb
--- /dev/null
+++ libgomp/testsuite/libgomp.c/target-1-foffload_disable.c
@@ -0,0 +1,3 @@
+/* { dg-additional-options "-foffload=disable" } */
+
+#include "target-1.c"
diff --git libgomp/testsuite/libgomp.c/target-foffload_disable.c
libgomp/testsuite/libgomp.c/target-foffload_disable.c
new file mode 100644
index 0000000..4a712da
--- /dev/null
+++ libgomp/testsuite/libgomp.c/target-foffload_disable.c
@@ -0,0 +1,18 @@
+/* { dg-additional-options "-foffload=disable" } */
+
+#include <omp.h>
+
+int main()
+{
+ if (!omp_is_initial_device())
+ __builtin_abort();
+#pragma omp target
+ {
+ if (!omp_is_initial_device())
+ __builtin_abort();
+ }
+ if (!omp_is_initial_device())
+ __builtin_abort();
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.fortran/target-foffload_disable.f
libgomp/testsuite/libgomp.fortran/target-foffload_disable.f
new file mode 100644
index 0000000..0d60534
--- /dev/null
+++ libgomp/testsuite/libgomp.fortran/target-foffload_disable.f
@@ -0,0 +1,14 @@
+! { dg-additional-options "-foffload=disable" }
+
+ PROGRAM MAIN
+ IMPLICIT NONE
+
+ INCLUDE "omp_lib.h"
+
+ IF (.NOT. OMP_IS_INITIAL_DEVICE()) CALL ABORT
+!$OMP TARGET
+ IF (.NOT. OMP_IS_INITIAL_DEVICE()) CALL ABORT
+!$OMP END TARGET
+ IF (.NOT. OMP_IS_INITIAL_DEVICE()) CALL ABORT
+
+ END
diff --git libgomp/testsuite/libgomp.fortran/target1-foffload_disable.f90
libgomp/testsuite/libgomp.fortran/target1-foffload_disable.f90
new file mode 100644
index 0000000..005328e
--- /dev/null
+++ libgomp/testsuite/libgomp.fortran/target1-foffload_disable.f90
@@ -0,0 +1,3 @@
+! { dg-additional-options "-cpp -foffload=disable" }
+
+#include "target1.f90"
diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp
libgomp/testsuite/libgomp.oacc-c++/c++.exp
index 88b0269..aa545a2 100644
--- libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -75,13 +75,12 @@ if { $lang_test_file_found } {
# Test OpenACC with available accelerators.
foreach offload_target_openacc $offload_targets_s_openacc {
- set tagopt "-DACC_DEVICE_TYPE_$offload_target_openacc=1"
-
- switch $offload_target_openacc {
- host {
+ switch -glob $offload_target_openacc {
+ disable {
set acc_mem_shared 1
+ set tagopt "-DACC_DEVICE_TYPE_host=1"
}
- nvidia {
+ nvptx* {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
# Don't bother; execution testing is going to FAIL.
untested "$subdir $offload_target_openacc offloading"
@@ -95,14 +94,13 @@ 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"
}
default {
set acc_mem_shared 0
}
}
- set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
-
- setenv ACC_DEVICE_TYPE $offload_target_openacc
+ set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared
-foffload=$offload_target_openacc"
dg-runtest $tests "$tagopt" "$libstdcxx_includes $DEFAULT_CFLAGS"
}
diff --git libgomp/testsuite/libgomp.oacc-c/c.exp
libgomp/testsuite/libgomp.oacc-c/c.exp
index 5020e6a..9d2065f 100644
--- libgomp/testsuite/libgomp.oacc-c/c.exp
+++ libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -38,13 +38,13 @@ set_ld_library_path_env_vars
set SAVE_ALWAYS_CFLAGS "$ALWAYS_CFLAGS"
foreach offload_target_openacc $offload_targets_s_openacc {
set ALWAYS_CFLAGS "$SAVE_ALWAYS_CFLAGS"
- set tagopt "-DACC_DEVICE_TYPE_$offload_target_openacc=1"
- switch $offload_target_openacc {
- host {
+ switch -glob $offload_target_openacc {
+ disable {
set acc_mem_shared 1
+ set tagopt "-DACC_DEVICE_TYPE_host=1"
}
- nvidia {
+ nvptx* {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
# Don't bother; execution testing is going to FAIL.
untested "$subdir $offload_target_openacc offloading"
@@ -58,14 +58,13 @@ 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"
}
default {
set acc_mem_shared 0
}
}
- set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
-
- setenv ACC_DEVICE_TYPE $offload_target_openacc
+ set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared
-foffload=$offload_target_openacc"
dg-runtest $tests "$tagopt" $DEFAULT_CFLAGS
}
diff --git libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index 2d6b647..3f678ba 100644
--- libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -67,13 +67,12 @@ if { $lang_test_file_found } {
# Test OpenACC with available accelerators.
foreach offload_target_openacc $offload_targets_s_openacc {
- set tagopt "-DACC_DEVICE_TYPE_$offload_target_openacc=1"
-
- switch $offload_target_openacc {
- host {
+ switch -glob $offload_target_openacc {
+ disable {
set acc_mem_shared 1
+ set tagopt "-DACC_DEVICE_TYPE_host=1"
}
- nvidia {
+ nvptx* {
if { ![check_effective_target_openacc_nvidia_accel_present] } {
# Don't bother; execution testing is going to FAIL.
untested "$subdir $offload_target_openacc offloading"
@@ -81,14 +80,13 @@ if { $lang_test_file_found } {
}
set acc_mem_shared 0
+ set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
}
default {
set acc_mem_shared 0
}
}
- set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
-
- setenv ACC_DEVICE_TYPE $offload_target_openacc
+ set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared
-foffload=$offload_target_openacc"
# For Fortran we're doing torture testing, as Fortran has far more tests
# with arrays etc. that testing just -O0 or -O2 is insufficient, that is
Grüße,
Thomas
signature.asc
Description: PGP signature
