This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |
| Other format: | [Raw text] | |
Hi!
Assuming that the overall approach (my option a) is fine, this is now
primarily a question about how to teach the driver to "the right thing".
(Joseph CCed as driver reviewer.)
On Wed, 5 Aug 2015 18:09:04 +0300, Ilya Verbin <iverbin@gmail.com> wrote:
> On Wed, Aug 05, 2015 at 10:40:44 +0200, Richard Biener wrote:
> > On Fri, Jul 31, 2015 at 4:20 PM, Ilya Verbin <iverbin@gmail.com> wrote:
> > > On Fri, Jul 31, 2015 at 16:08:27 +0200, Thomas Schwinge wrote:
> > >> We had established the use of a boolean flag have_offload in gcc::context
> > >> to indicate whether during compilation, we've actually seen any code to
> > >> be offloaded (see cited below the relevant parts of the patch by Ilya et
> > >> al.). This means that currently, the whole offload machinery will not be
> > >> run unless we actually have any offloaded data. This means that the
> > >> configured mkoffload programs (-foffload=[...], defaulting to
> > >> configure-time --enable-offload-targets=[...]) will not be invoked unless
> > >> we actually have any offloaded data. This means that we will not
> > >> actually generate constructor code to call libgomp's
> > >> GOMP_offload_register unless we actually have any offloaded data.
> > >
> > > Yes, that was the plan.
> > >
> > >> runtime, in libgomp, we then cannot reliably tell which -foffload=[...]
> > >> targets have been specified during compilation.
> > >>
> > >> But: at runtime, I'd like to know which -foffload=[...] targets have been
> > >> specified during compilation, so that we can, for example, reliably
> > >> resort to host fallback execution for -foffload=disable instead of
> > >> getting error message that an offloaded function is missing.
> > >
> > > It's easy to fix:
> > >
> > > diff --git a/libgomp/target.c b/libgomp/target.c
> > > index a5fb164..f81d570 100644
> > > --- a/libgomp/target.c
> > > +++ b/libgomp/target.c
> > > @@ -1066,9 +1066,6 @@ gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
> > > k.host_end = k.host_start + 1;
> > > splay_tree_key tgt_fn = splay_tree_lookup (&devicep->mem_map, &k);
> > > gomp_mutex_unlock (&devicep->lock);
> > > - if (tgt_fn == NULL)
> > > - gomp_fatal ("Target function wasn't mapped");
> > > -
> > > return (void *) tgt_fn->tgt_offset;
> > > }
Won't that possibly result in a NULL pointer dereference (tgt_fn) --
instead return NULL, I think?
> > > @@ -1095,6 +1092,8 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
> > > return gomp_target_fallback (fn, hostaddrs);
> > >
> > > void *fn_addr = gomp_get_target_fn_addr (devicep, fn);
> > > + if (fn_addr == NULL)
> > > + return gomp_target_fallback (fn, hostaddrs);
Is that reliable? Consider the following scenario, with f1 and f2
implemented in separate TUs:
#pragma omp target data [map clauses]
{
f1([...]);
f2([...]);
}
Consider that in f1 we have a OpenMP target region with offloading data
available, and in f2 we have a OpenMP target region without offloading
data available. In this case, the GOMP_target in f1 will execute on the
offloading target, but the GOMP_target in f2 will resort to host fallback
-- and we then likely have data inconsistencies, as the data specified by
the map clauses is not synchronized between host and device.
Admittedly, this is user error (inconsistent set of offloading functions
available -- need either all, or none), but in such a scenario probably
we should be doing a better job (at detecting this). (Note, I'm not sure
whether my current patch actually does any better.) ;-)
> > >> other hand, for example, for -foffload=nvptx-none, even if user program
> > >> code doesn't contain any offloaded data (and thus the offload machinery
> > >> has not been run), the user program might still contain any executable
> > >> directives or OpenACC runtime library calls, so we'd still like to use
> > >> the libgomp nvptx plugin. However, we currently cannot detect this
> > >> situation.
> > >>
> > >> I see two ways to resolve this: a) embed the compile-time -foffload=[...]
> > >> configuration in the executable (as a string, for example) for libgomp to
> > >> look that up, or b) make it a requirement that (if configured via
> > >> -foffload=[...]), the offload machinery is run even if there is not
> > >> actually any data to be offloaded, so we then reliably get the respective
> > >> constructor call to libgomp's GOMP_offload_register. I once began to
> > >> implement a), but this to get a big ugly, so then looked into b) instead.
> > >> Compared to the status quo, always running the whole offloading machinery
> > >> for the configured -foffload=[...] targets whenever -fopenacc/-fopenmp
> > >> are active, certainly does introduce some overhead when there isn't
> > >> actually any code to be offloaded, so I'm not sure whether that is
> > >> acceptable?
> > >
> > > I vote for (a).
OK. Any other opinions?
> > What happens for conflicting -fofffload=[...] options in different TUs?
>
> If you're asking about what happens now, only the list of offload targets from
> link-time -foffload=tgt1,tgt2 option matters.
I'm fine with that -- require the user to specify a consistent set of
-foffload options. (Consistent in the sense such that all offload data
must be available that can possibly be required at run-time.)
> I don't like plan (b) because it calls ipa_write_summaries unconditionally for
> all OpenMP programs, which creates IR sections, which increases filesize and may
> cause other problems, e.g. <https://gcc.gnu.org/bugzilla/show_bug.cgi?id=63868>.
> Also compile-time is increased because of LTO machinery, mkoffloads, etc.
I think the compile-time effort is the only real argument against this --
everything else are "just" bugs that need to/have been addressed. But
OK, assuming it is doable, I have a preference for option a), too.
(Which is why it's option a).) ;-)
> If OpenACC requires some registration in libgomp even without offload, maybe you
> can run this machinery only under flag_openacc?
That seems to much special-casing for my taste.
Here is an attempt at option a). The idea is to have "the compilation
process" create a constructor call to a (new) libgomp
GOMP_set_offload_targets function, and at run-time, the information
passed with that call will be used to determine the set of requested
offload targets (if such a GOMP_set_offload_targets call is not made,
defaulting to all of them, just like now).
I've settled on adding a function combined with a constructor call, as
that seems the easiest solution without adding a lot of complexity, to
not break the ABI for older (pre-GCC 6) executables dynamically linking
against the new (GCC 6) libgomp. (If that weren't a concern, we could
have embedded a string GOMP_offload_targets in the final executable, and
have libgomp look up the information from there, but that string won't be
available in pre-GCC 6 executables, which could possibly be deal with by
relying on symbol lookup ordering (provide a default for
GOMP_offload_targets in libgomp itself -- fragile), or using weak
symbols, but those "are difficult"/not generally supported...)
A lot of the libgomp changes (which apply to gomp-4_0-branch) are just
cleanup, so don't be scared by that. (That is, the trunk patch will be
simpler in that regard.)
The interesting thing, that I'd like comments on, are the driver changes.
I'm adding a new add-omp-infile spec function (add_omp_infile_spec_func),
which is "invoked" whenever we're linking against libgomp (-fopenacc or
-fopenmp or -ftree-parallelize-loops). (Should probably generally clean
that up a little.)
This function Âgenerate[s] a C source file containing a constructor call
to GOMP_set_offload_targets [...], and adds that as an infileÂ. This
"basically" works ;-) -- but really only for C source code, and for C++
and Fortran it fails if there are command-line options used that conflict
with the C compilation of add-omp-infile, such as (from a libgomp
testsuite run): for C++: -std=c++11, -fno-extern-tls-init, or for
Fortran: -fcray-pointer, -fintrinsic-modules-path. Any suggestion about
how to overcome that?
That is, in the driver, how can I compile this C file with just the
options that are applicable to the C compiler, and then just link in the
object file?
gcc/gcc.c | 104 +++++++++++--
libgomp/config.h.in | 8 +-
libgomp/configure | 33 +++-
libgomp/env.c | 6 +-
libgomp/libgomp.h | 1 +
libgomp/libgomp.map | 7 +-
libgomp/libgomp_g.h | 1 +
libgomp/oacc-init.c | 18 ++-
libgomp/plugin/configfrag.ac | 10 +-
libgomp/target.c | 172 ++++++++++++++++-----
libgomp/testsuite/lib/libgomp.exp | 75 ++-------
libgomp/testsuite/libgomp.c++/c++.exp | 13 --
libgomp/testsuite/libgomp.c/c.exp | 2 -
libgomp/testsuite/libgomp.fortran/fortran.exp | 5 -
libgomp/testsuite/libgomp.graphite/graphite.exp | 2 -
libgomp/testsuite/libgomp.oacc-c++/c++.exp | 33 ++--
libgomp/testsuite/libgomp.oacc-c/c.exp | 17 +-
libgomp/testsuite/libgomp.oacc-fortran/fortran.exp | 23 +--
18 files changed, 322 insertions(+), 208 deletions(-)
diff --git gcc/gcc.c gcc/gcc.c
index 0642be1..bcb32f5 100644
--- gcc/gcc.c
+++ gcc/gcc.c
@@ -1,3 +1,5 @@
+#define FPRINTF if (getenv("DEBUG")) fprintf
+
/* Compiler driver program that can handle many languages.
Copyright (C) 1987-2015 Free Software Foundation, Inc.
@@ -158,7 +160,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION;
static const char *spec_machine = DEFAULT_TARGET_MACHINE;
static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE;
-/* List of offload targets. */
+/* List of offload targets. Empty string for -foffload=disable. */
static char *offload_targets = NULL;
@@ -275,6 +277,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
@@ -1061,6 +1065,14 @@ 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 the user didn't specify any, default to all configured offload
+ targets. */
+ "%{!foffload=*:-foffload=" OFFLOAD_TARGETS "}",
+ /* 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
};
@@ -1491,6 +1503,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
@@ -3450,6 +3463,8 @@ static int last_language_n_infiles;
static void
handle_foffload_option (const char *arg)
{
+ FPRINTF(stderr, "%s (\"%s\")\n", __FUNCTION__, arg);
+
const char *c, *cur, *n, *next, *end;
char *target;
@@ -7162,6 +7177,8 @@ driver::build_multilib_strings () const
void
driver::set_up_specs () const
{
+ FPRINTF(stderr, "%s\n", __FUNCTION__);
+
const char *spec_machine_suffix;
char *specs_file;
size_t i;
@@ -7448,22 +7465,16 @@ driver::maybe_putenv_COLLECT_LTO_WRAPPER () const
void
driver::maybe_putenv_OFFLOAD_TARGETS () const
{
- const char *targets = offload_targets;
-
- /* If no targets specified by -foffload, use all available targets. */
- if (!targets)
- targets = OFFLOAD_TARGETS;
+ FPRINTF(stderr, "OFFLOAD_TARGETS = %s\n", offload_targets);
- if (strlen (targets) > 0)
+ if (offload_targets && offload_targets[0] != '\0')
{
obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
- obstack_grow (&collect_obstack, targets,
- strlen (targets) + 1);
+ obstack_grow (&collect_obstack, offload_targets,
+ strlen (offload_targets) + 1);
xputenv (XOBFINISH (&collect_obstack, char *));
}
-
- free (offload_targets);
}
/* Reject switches that no pass was interested in. */
@@ -9507,6 +9518,77 @@ greater_than_spec_func (int argc, const char **argv)
return NULL;
}
+/* If applicable, generate a C source file containing a constructor call to
+ GOMP_set_offload_targets, to inform libgomp which offload targets have
+ actually been requested (-foffload=[...]), and adds 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, !save_temps_flag);
+ 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_set_offload_targets (const char *);\n"
+ "static __attribute__ ((constructor)) void\n"
+ "init (void)\n"
+ "{\n"
+ " GOMP_set_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);
+
+ //TODO: correct thing to do?
+#if 1
+ add_infile (tmp_filename, "cpp-output");
+ return NULL;
+#elif 0
+ return tmp_filename;
+#elif 0
+ store_arg ("-x", 0, 0);
+ store_arg ("cpp-output", 0, 0);
+ store_arg (tmp_filename, 0, 0);
+ return NULL;
+#else
+ //add_infile (tmp_filename, /* TODO */ "cpp-output");
+ //int i = n_infiles - 1;
+ //input_file_number = i;
+ set_input (tmp_filename);
+ //outfiles[i] = gcc_input_filename;
+ input_file_compiler
+ = lookup_compiler (gcc_input_filename, input_filename_length,
+ "cpp-output");
+ err = do_spec (input_file_compiler->spec);
+ //infiles[i].compiled = true;
+ if (err < 0)
+ {
+ delete_failure_queue ();
+ errorcount++;
+ }
+ clear_failure_queue ();
+#endif
+}
+
/* Insert backslash before spaces in ORIG (usually a file path), to
avoid being broken by spec parser.
diff --git libgomp/config.h.in libgomp/config.h.in
index 8533f03..d9d5914 100644
--- libgomp/config.h.in
+++ libgomp/config.h.in
@@ -24,6 +24,12 @@
/* Define to 1 if you have the <dlfcn.h> header file. */
#undef HAVE_DLFCN_H
+/* Define to 1 if you have the `fnmatch' function. */
+#undef HAVE_FNMATCH
+
+/* Define to 1 if you have the <fnmatch.h> header file. */
+#undef HAVE_FNMATCH_H
+
/* Define to 1 if you have the `getloadavg' function. */
#undef HAVE_GETLOADAVG
@@ -95,7 +101,7 @@
*/
#undef LT_OBJDIR
-/* Define to hold the list of target names suitable for offloading. */
+/* Define to hold the list of offload targets, separated by colons. */
#undef OFFLOAD_TARGETS
/* Name of package */
diff --git libgomp/configure libgomp/configure
index c93e877..3d990bb 100755
--- libgomp/configure
+++ libgomp/configure
@@ -15119,6 +15119,33 @@ esac
offload_targets=
plugin_support=yes
+for ac_header in fnmatch.h
+do :
+ ac_fn_c_check_header_mongrel "$LINENO" "fnmatch.h" "ac_cv_header_fnmatch_h" "$ac_includes_default"
+if test "x$ac_cv_header_fnmatch_h" = x""yes; then :
+ cat >>confdefs.h <<_ACEOF
+#define HAVE_FNMATCH_H 1
+_ACEOF
+
+else
+ plugin_support=no
+fi
+
+done
+
+for ac_func in fnmatch
+do :
+ ac_fn_c_check_func "$LINENO" "fnmatch" "ac_cv_func_fnmatch"
+if test "x$ac_cv_func_fnmatch" = x""yes; then :
+ cat >>confdefs.h <<_ACEOF
+#define HAVE_FNMATCH 1
+_ACEOF
+
+else
+ plugin_support=no
+fi
+done
+
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for dlsym in -ldl" >&5
$as_echo_n "checking for dlsym in -ldl... " >&6; }
if test "${ac_cv_lib_dl_dlsym+set}" = set; then :
@@ -15236,10 +15263,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 +15307,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/env.c libgomp/env.c
index 1811bf5..6b5e963 100644
--- libgomp/env.c
+++ libgomp/env.c
@@ -1175,11 +1175,7 @@ handle_omp_display_env (unsigned long stacksize, int wait_policy)
}
-/* TODO. See testsuite/lib/libgomp.exp:libgomp_init. */
-#if 0
-static
-#endif
-void __attribute__((constructor))
+static void __attribute__((constructor))
initialize_env (void)
{
unsigned long thread_limit_var, stacksize;
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 420a525..1f95906 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -785,6 +785,7 @@ extern void gomp_unmap_vars (struct target_mem_desc *, bool);
extern void gomp_init_device (struct gomp_device_descr *);
extern void gomp_fini_device (struct gomp_device_descr *);
extern void gomp_unload_device (struct gomp_device_descr *);
+extern bool gomp_offload_target_available_p (int);
/* work.c */
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 36932d3..05e75fb0 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -339,6 +339,7 @@ GOACC_2.0.GOMP_4_BRANCH {
GOACC_get_ganglocal_ptr;
GOACC_parallel_keyed;
GOACC_register_static;
+ GOMP_set_offload_targets;
} GOACC_2.0;
GOMP_PLUGIN_1.0 {
@@ -352,9 +353,3 @@ GOMP_PLUGIN_1.0 {
GOMP_PLUGIN_async_unmap_vars;
GOMP_PLUGIN_acc_thread;
};
-
-# TODO. See testsuite/lib/libgomp.exp:libgomp_init.
-INTERNAL {
- global:
- initialize_env;
-};
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index e65b888..e67fc86 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_set_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 8dfe4a7..78f130c 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_available_p. (That is, hard
+ error if not actually available.) */
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_available_p, to not decide on an offload
+ target that we don't have offload data available for. */
+ && gomp_offload_target_available_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 8c2a420..e2392e1 100644
--- libgomp/plugin/configfrag.ac
+++ libgomp/plugin/configfrag.ac
@@ -29,6 +29,8 @@
offload_targets=
AC_SUBST(offload_targets)
plugin_support=yes
+AC_CHECK_HEADERS([fnmatch.h], , [plugin_support=no])
+AC_CHECK_FUNCS([fnmatch], , [plugin_support=no])
AC_CHECK_LIB(dl, dlsym, , [plugin_support=no])
if test x"$plugin_support" = xyes; then
AC_DEFINE(PLUGIN_SUPPORT, 1,
@@ -92,10 +94,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 +127,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 +141,7 @@ if test x"$enable_offload_targets" != x; then
done
fi
AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets",
- [Define to hold the list of target names suitable for offloading.])
+ [Define to hold the list of 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 6426254..9cf5251 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -41,6 +41,7 @@
#ifdef PLUGIN_SUPPORT
#include <dlfcn.h>
+#include <fnmatch.h>
#include "plugin-suffix.h"
#endif
@@ -122,17 +123,26 @@ 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;
+ /* If the device-var ICV does not actually have offload data available, don't
+ try use it (which will fail), and use host fallback instead. */
+ if (device == GOMP_DEVICE_ICV
+ && !gomp_offload_target_available_p (devices[device_id].type))
+ return NULL;
+
return &devices[device_id];
}
@@ -947,6 +957,49 @@ gomp_fini_device (struct gomp_device_descr *devicep)
devicep->is_initialized = false;
}
+/* Do we have offload data available for the given offload target type?
+ Instead of verifying that *all* offload data is available that could
+ possibly be required, we instead just look for *any*. If we later find any
+ offload data missing, that's user error. */
+
+attribute_hidden bool
+gomp_offload_target_available_p (int type)
+{
+ bool available = false;
+
+ /* Has the offload target already been initialized? */
+ for (int i = 0; !available && i < num_devices; i++)
+ {
+ struct gomp_device_descr *devicep = &devices[i];
+ gomp_mutex_lock (&devicep->lock);
+ if (devicep->type == type && devicep->is_initialized)
+ available = true;
+ gomp_mutex_unlock (&devicep->lock);
+ }
+
+ if (!available)
+ {
+ // TODO: locking correct?
+ gomp_mutex_lock (®ister_lock);
+
+ /* If there is no offload data available at all, we cannot later fail to
+ find any of it for a specific offload target. This is the case where
+ there are no offloaded code regions in user code, but there can still
+ be executable directives used, or runtime library calls made. */
+ if (num_offload_images == 0)
+ available = true;
+
+ /* Can the offload target be initialized? */
+ for (int i = 0; !available && i < num_offload_images; i++)
+ if (offload_images[i].type == type)
+ available = true;
+
+ gomp_mutex_unlock (®ister_lock);
+ }
+
+ return available;
+}
+
/* 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
@@ -1116,6 +1169,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);
@@ -1212,6 +1267,38 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
return 0;
}
+/* Helper, to translate from an offload target to the corresponding plugin name. */
+
+static const char *
+offload_target_to_plugin_name (const char *offload_target)
+{
+ if (fnmatch ("*-intelmic*", offload_target, 0) == 0)
+ return "intelmic";
+ if (fnmatch ("nvptx*", offload_target, 0) == 0)
+ return "nvptx";
+ gomp_fatal ("Unknown offload target: %s", offload_target);
+}
+
+/* List of offload targets, separated by colon. Defaults to the list
+ determined when configuring libgomp. */
+static const char *gomp_offload_targets = OFFLOAD_TARGETS;
+
+/* 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);
+
+ //TODO: any locking?
+ /* Make sure this gets called early. */
+ assert (gomp_is_initialized == PTHREAD_ONCE_INIT);
+ ///* Make sure this only gets called once. */
+ //assert (gomp_offload_targets == OFFLOAD_TARGETS);
+ gomp_offload_targets = offload_targets;
+}
+
/* 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
@@ -1219,11 +1306,12 @@ 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;
@@ -1231,48 +1319,60 @@ gomp_target_init (void)
num_devices = 0;
devices = NULL;
- cur = OFFLOAD_TARGETS;
+ cur = gomp_offload_targets;
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 = cur + strlen (cur);
+
+ 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);
+ plugin_name = gomp_malloc (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... */
+ const char *cur_plugin_name
+ = offload_target_to_plugin_name (plugin_name
+ + gomp_plugin_prefix_len);
+ 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 (devices,
+ ((num_devices + new_num_devices)
+ * sizeof (struct gomp_device_descr)));
for (i = 0; i < new_num_devices; i++)
{
current_device.target_id = i;
@@ -1286,18 +1386,12 @@ 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_malloc (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 33d1a54..898dfc3 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* {
# TODO. 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
@@ -235,56 +232,6 @@ proc libgomp_init { args } {
if { $offload_additional_options != "" } {
lappend ALWAYS_CFLAGS "additional_flags=${offload_additional_options}"
}
-
- # TODO. Evil hack. DejaGnu doesn't have a mechanism for setting
- # environment variables on remote boards. Thus, we have to fake it, using
- # GCC's constructor attributes to create object files that install the
- # desired environment variables.
- set e_list [list \
- [list defaults DUMMY=dummy ] \
- [list ACC_DEVICE_TYPE-host ACC_DEVICE_TYPE=host ] \
- [list ACC_DEVICE_TYPE-nvidia ACC_DEVICE_TYPE=nvidia ] ]
- foreach e $e_list {
- set v [lindex $e 0]
- set s [lindex $e 1]
- verbose "creating constructor-setenv: $v: $s"
- set src constructor-setenv-$v.c
- set obj constructor-setenv-$v.o
- set f_src [open $src "w"]
- puts $f_src "static void __attribute__((constructor(1000)))"
- puts $f_src "init_env(void) {"
- puts $f_src " int putenv(char *);"
- puts $f_src " putenv(\"$s\");"
- puts $f_src "}"
- if { $v == "defaults" } {
- # TODO. We want libgomp to initialize after the putenv calls.
- # But: shared libraries' constructors (and thus
- # env.c:initialize_env) will be called before the executable's
- # (init_env functions created above), so it will already have been
- # initialized (and has to be, in case we're not linking in this
- # gunk). Assuming no execution of other libgomp functionality in
- # between (which we're not doing during initialization),
- # initialize_env's effects are idempotent when calling it again, so
- # we'll do that now, after the putenv calls have been executed.
- puts $f_src "static void __attribute__((constructor(1001)))"
- puts $f_src "init_libgomp(void) {"
- # Some test cases specify -fno-openmp, so libgomp isn't linked in.
- puts $f_src " void initialize_env(void) __attribute__((weak));"
- puts $f_src " if (initialize_env)"
- puts $f_src " initialize_env();"
- puts $f_src "}"
- }
- close $f_src
- # TODO. Using whichever compiler is currently configured... At least
- # switch it into C mode.
- set lines [libgomp_target_compile $src $obj object "additional_flags=-xc"]
- # TODO. Error checking.
- file delete $src
- }
- # When adding constructor-setenv-*.o files, make sure to cancel any -x flag
- # that may have been set before.
- lappend ALWAYS_CFLAGS "ldflags=-x none"
- lappend ALWAYS_CFLAGS "ldflags=constructor-setenv-defaults.o"
}
#
@@ -296,6 +243,7 @@ proc libgomp_target_compile { source dest type options } {
global libgomp_compile_options
global gluefile wrap_flags
global ALWAYS_CFLAGS
+ global GCC_UNDER_TEST
global lang_test_file
global lang_library_path
global lang_link_flags
@@ -323,6 +271,7 @@ proc libgomp_target_compile { source dest type options } {
lappend options "additional_flags=[libio_include_flags]"
lappend options "timeout=[timeout_value]"
+ lappend options "compiler=$GCC_UNDER_TEST"
set options [concat $libgomp_compile_options $options]
@@ -370,7 +319,7 @@ proc check_effective_target_offload_device { } {
proc check_effective_target_openacc_nvidia_accel_supported { } {
global offload_targets_s_openacc
- set res [lsearch $offload_targets_s_openacc "nvidia" ]
+ set res [lsearch -glob $offload_targets_s_openacc "nvptx*" ]
if { $res != -1 } {
return 1;
}
@@ -396,7 +345,7 @@ proc check_effective_target_openacc_nvidia_accel_selected { } {
return 0;
}
global offload_target_openacc
- if { $offload_target_openacc == "nvidia" } {
+ if { [string match "nvptx*" $offload_target_openacc] } {
return 1;
}
return 0;
@@ -406,7 +355,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++/c++.exp libgomp/testsuite/libgomp.c++/c++.exp
index d6d525a..0454f95 100644
--- libgomp/testsuite/libgomp.c++/c++.exp
+++ libgomp/testsuite/libgomp.c++/c++.exp
@@ -4,7 +4,6 @@ load_gcc_lib gcc-dg.exp
global shlib_ext
set shlib_ext [get_shlib_extension]
-#TODO
set lang_link_flags "-lstdc++"
set lang_test_file_found 0
set lang_library_path "../libstdc++-v3/src/.libs"
@@ -47,13 +46,6 @@ if { $blddir != "" } {
}
if { $lang_test_file_found } {
- if ![info exists GXX_UNDER_TEST] then {
- # TODO. See libgomp.oacc-c++/c++.exp.
- set HAVE_SET_GXX_UNDER_TEST ""
- set GXX_UNDER_TEST "$GCC_UNDER_TEST"
- }
- lappend libgomp_compile_options "compiler=$GXX_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [find $srcdir/$subdir *.C]]
@@ -76,10 +68,5 @@ if { $lang_test_file_found } {
dg-runtest $tests "" "$libstdcxx_includes $DEFAULT_CFLAGS"
}
-# TODO. See above.
-if { [info exists HAVE_SET_GXX_UNDER_TEST] } {
- unset GXX_UNDER_TEST
-}
-
# All done.
dg-finish
diff --git libgomp/testsuite/libgomp.c/c.exp libgomp/testsuite/libgomp.c/c.exp
index 25f347b..300b921 100644
--- libgomp/testsuite/libgomp.c/c.exp
+++ libgomp/testsuite/libgomp.c/c.exp
@@ -23,8 +23,6 @@ dg-init
# Turn on OpenMP.
lappend ALWAYS_CFLAGS "additional_flags=-fopenmp"
-lappend libgomp_compile_options "compiler=$GCC_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [find $srcdir/$subdir *.c]]
diff --git libgomp/testsuite/libgomp.fortran/fortran.exp libgomp/testsuite/libgomp.fortran/fortran.exp
index 883c416..f684abc 100644
--- libgomp/testsuite/libgomp.fortran/fortran.exp
+++ libgomp/testsuite/libgomp.fortran/fortran.exp
@@ -47,11 +47,6 @@ if { $blddir != "" } {
}
if { $lang_test_file_found } {
- if ![info exists GFORTRAN_UNDER_TEST] then {
- set GFORTRAN_UNDER_TEST $GCC_UNDER_TEST
- }
- lappend libgomp_compile_options "compiler=$GFORTRAN_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [find $srcdir/$subdir *.\[fF\]{,90,95,03,08}]]
diff --git libgomp/testsuite/libgomp.graphite/graphite.exp libgomp/testsuite/libgomp.graphite/graphite.exp
index 716cdc3..d737c85 100644
--- libgomp/testsuite/libgomp.graphite/graphite.exp
+++ libgomp/testsuite/libgomp.graphite/graphite.exp
@@ -48,8 +48,6 @@ dg-init
# Turn on OpenMP.
lappend ALWAYS_CFLAGS "additional_flags=-fopenmp"
-lappend libgomp_compile_options "compiler=$GCC_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [find $srcdir/$subdir *.c]]
diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp libgomp/testsuite/libgomp.oacc-c++/c++.exp
index e5c875c..f513d87 100644
--- libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -13,7 +13,6 @@ load_gcc_lib gcc-dg.exp
global shlib_ext
set shlib_ext [get_shlib_extension]
-#TODO
set lang_link_flags "-lstdc++"
set lang_test_file_found 0
set lang_library_path "../libstdc++-v3/src/.libs"
@@ -32,6 +31,11 @@ dg-init
# Turn on OpenACC.
lappend ALWAYS_CFLAGS "additional_flags=-fopenacc"
+# Switch into C++ mode. Otherwise, the libgomp.oacc-c-c++-common/*.c
+# files would be compiled as C files.
+set SAVE_GCC_UNDER_TEST "$GCC_UNDER_TEST"
+set GCC_UNDER_TEST "$GCC_UNDER_TEST -x c++"
+
set blddir [lookfor_file [get_multilibs] libgomp]
@@ -56,14 +60,6 @@ if { $blddir != "" } {
}
if { $lang_test_file_found } {
- if ![info exists GXX_UNDER_TEST] then {
- # Use GCC_UNDER_TEST, but switch into C++ mode, as otherwise the
- # libgomp.oacc-c-c++-common/*.c files would be compiled as C files.
- set HAVE_SET_GXX_UNDER_TEST ""
- set GXX_UNDER_TEST "$GCC_UNDER_TEST -x c++"
- }
- lappend libgomp_compile_options "compiler=$GXX_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [concat \
[find $srcdir/$subdir *.C] \
@@ -104,17 +100,14 @@ if { $lang_test_file_found } {
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"
- # Set $ACC_DEVICE_TYPE. See the comments in
- # ../lib/libgomp.exp:libgomp_init.
- lappend ALWAYS_CFLAGS "ldflags=constructor-setenv-ACC_DEVICE_TYPE-$offload_target_openacc.o"
# Todo: Determine shared memory or not using run-time test.
- 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"
@@ -128,12 +121,14 @@ 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
+ #TODO error
}
}
- set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
+ set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared -foffload=$offload_target_openacc"
dg-runtest $tests "$tagopt" "$libstdcxx_includes $DEFAULT_CFLAGS"
gcc-dg-runtest $ttests "$tagopt" "$libstdcxx_includes"
@@ -141,9 +136,7 @@ if { $lang_test_file_found } {
}
# See above.
-if { [info exists HAVE_SET_GXX_UNDER_TEST] } {
- unset GXX_UNDER_TEST
-}
+set GCC_UNDER_TEST "$SAVE_GCC_UNDER_TEST"
unset TORTURE_OPTIONS
diff --git libgomp/testsuite/libgomp.oacc-c/c.exp libgomp/testsuite/libgomp.oacc-c/c.exp
index c91a41b..03fe3a4 100644
--- libgomp/testsuite/libgomp.oacc-c/c.exp
+++ libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -32,8 +32,6 @@ dg-init
# Turn on OpenACC.
lappend ALWAYS_CFLAGS "additional_flags=-fopenacc"
-lappend libgomp_compile_options "compiler=$GCC_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [concat \
[find $srcdir/$subdir *.c] \
@@ -62,17 +60,14 @@ 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"
- # Set $ACC_DEVICE_TYPE. See the comments in
- # ../lib/libgomp.exp:libgomp_init.
- lappend ALWAYS_CFLAGS "ldflags=constructor-setenv-ACC_DEVICE_TYPE-$offload_target_openacc.o"
# Todo: Determine shared memory or not using run-time test.
- 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"
@@ -86,12 +81,14 @@ 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
+ #TODO error
}
}
- set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
+ set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared -foffload=$offload_target_openacc"
dg-runtest $tests "$tagopt" $DEFAULT_CFLAGS
gcc-dg-runtest $ttests "$tagopt" ""
diff --git libgomp/testsuite/libgomp.oacc-fortran/fortran.exp libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index df46004..b9ad6dc 100644
--- libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -49,11 +49,6 @@ if { $blddir != "" } {
}
if { $lang_test_file_found } {
- if ![info exists GFORTRAN_UNDER_TEST] then {
- set GFORTRAN_UNDER_TEST $GCC_UNDER_TEST
- }
- lappend libgomp_compile_options "compiler=$GFORTRAN_UNDER_TEST"
-
# Gather a list of all tests.
set tests [lsort [find $srcdir/$subdir *.\[fF\]{,90,95,03,08}]]
@@ -74,20 +69,14 @@ if { $lang_test_file_found } {
set_ld_library_path_env_vars
# Test OpenACC with available accelerators.
- 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"
- # Set $ACC_DEVICE_TYPE. See the comments in
- # ../lib/libgomp.exp:libgomp_init.
- lappend ALWAYS_CFLAGS "ldflags=constructor-setenv-ACC_DEVICE_TYPE-$offload_target_openacc.o"
-
# Todo: Determine shared memory or not using run-time test.
- 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,12 +84,14 @@ if { $lang_test_file_found } {
}
set acc_mem_shared 0
+ set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
}
default {
set acc_mem_shared 0
+ #TODO error
}
}
- set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
+ 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
Attachment:
pgpKXTtWjbhML.pgp
Description: PGP signature
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |