From: Julian Brown Date: Thu, 28 Feb 2019 16:53:54 +0000 (-0800) Subject: [PATCH] Forward -foffload=[...] from the driver (compile-time) to libgomp (run-time) X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=a8b69300ec4a79512d9407c862967f2fcbac659e;p=thirdparty%2Fgcc.git [PATCH] Forward -foffload=[...] from the driver (compile-time) to libgomp (run-time) gcc/ PR libgomp/81886 * doc/invoke.texi (-ffixed-@var{reg}): Document conflict with Fortran options. * gcc.c (add_omp_infile_spec_func, spec_lang_mask_accept): New. (driver_self_specs): Add spec to use %:add-omp-infile(). (static_spec_functions): Add add-omp-infile. (struct switchstr): Add lang_mask field. Expand comment. (struct infile): Add lang_mask field. (add_infile, save_switch, do_spec): Add lang_mask argument. (driver_unknown_option_callback, driver_wrong_lang_callback) (driver_handle_option, process_command, do_self_spec) (driver::do_spec_on_infiles, driver::maybe_run_linker): All callers changed. (give_switch): Check languages of switch against spec_lang_mask_accept. (driver::maybe_putenv_OFFLOAD_TARGETS): Don't free offload_targets. * gcc.h (do_spec): Update prototype. gcc/fortran/ PR libgomp/81886 * gfortranspec.c (lang_specific_pre_link): Update call to do_spec. libgomp/ PR libgomp/81886 * plugin/configfrag.ac (tgt_name): Do not set. (offload_targets): Separate with colons not commas. * config.h.in, configure: Regenerate. * libgomp.h (gomp_offload_target_available_p): New prototype. * libgomp.map (GOACC_2.0.GOMP_4_BRANCH): Add GOMP_set_offload_targets. * libgomp_g.h (GOMP_set_offload_targets): New prototype. * oacc-init.c (resolve_device): Use gomp_offload_target_available_p. * target.c (resolve_device): Use host fallback when offload data not available. (gomp_offload_target_available_p, offload_target_to_plugin_name) (gomp_offload_targets, gomp_offload_targets_init) (GOMP_set_offload_targets, gomp_plugin_prefix) (gomp_plugin_suffix): New. (gomp_load_plugin_for_device): Add gomp_debug call. (gomp_target_init): Use gomp_offload_targets instead of OFFLOAD_TARGETS. Handle and rewrie colon-separated string. * 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. * testsuite/lib/libgomp.exp: Expect offload targets to be colon-separated. Adjust matching of offload targets. (libgomp_init) (check_effective_target_openacc_nvidia_accel_configured) (check_effective_target_openacc_host_selected): Adjust checks of offload target names. * testsuite/libgomp.oacc-c++/c++.exp: Adjust set of offload targets. Use -foffload instead of setenv ACC_DEVICE_TYPE. 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-c++-common/parallel-dims.c: Adapt. * 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-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. (cherry picked from gomp-4_0-branch r227045, r227154, r227175, r248030, and openacc-gcc-7-branch commit 917e247055a37f912129ed545719182de0046adb) (cherry picked from openacc-gcc-9-branch commit 789c1d022a871eb06ab08bbb63dcb89006361d93) --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp new file mode 100644 index 000000000000..712ac488baf4 --- /dev/null +++ b/gcc/ChangeLog.omp @@ -0,0 +1,22 @@ +2015-08-20 Thomas Schwinge + Joseph Myers + + PR libgomp/81886 + * doc/invoke.texi (-ffixed-@var{reg}): Document conflict with + Fortran options. + * gcc.c (add_omp_infile_spec_func, spec_lang_mask_accept): New. + (driver_self_specs): Add spec to use %:add-omp-infile(). + (static_spec_functions): Add add-omp-infile. + (struct switchstr): Add lang_mask field. Expand comment. + (struct infile): Add lang_mask field. + (add_infile, save_switch, do_spec): Add lang_mask argument. + (driver_unknown_option_callback, driver_wrong_lang_callback) + (driver_handle_option, process_command, do_self_spec) + (driver::do_spec_on_infiles, driver::maybe_run_linker): All + callers changed. + (give_switch): Check languages of switch against + spec_lang_mask_accept. + (driver::maybe_putenv_OFFLOAD_TARGETS): Don't free + offload_targets. + * gcc.h (do_spec): Update prototype. + diff --git a/gcc/d/d-spec.cc b/gcc/d/d-spec.cc index 9eba6902bb9f..e423852eae6a 100644 --- a/gcc/d/d-spec.cc +++ b/gcc/d/d-spec.cc @@ -505,7 +505,7 @@ int lang_specific_pre_link (void) { if ((phobos_library != PHOBOS_NOLINK && need_phobos) || need_spec) - do_spec ("%:include(libgphobos.spec)"); + do_spec ("%:include(libgphobos.spec)", 0); return 0; } diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index 29585cf15aac..08fab1f80e57 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -14089,6 +14089,10 @@ macro in the machine description macro file. This flag does not have a negative form, because it specifies a three-way choice. +Note that this flag may conflict with the @option{-ffixed-form} as +well as @option{-ffixed-line-length-none} and +@option{-ffixed-line-length-} options of the Fortran front end. + @item -fcall-used-@var{reg} @opindex fcall-used Treat the register named @var{reg} as an allocable register that is diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp new file mode 100644 index 000000000000..88dfde119f13 --- /dev/null +++ b/gcc/fortran/ChangeLog.omp @@ -0,0 +1,5 @@ +2015-08-20 Joseph Myers + + PR libgomp/81886 + * gfortranspec.c (lang_specific_pre_link): Update call to do_spec. + diff --git a/gcc/fortran/gfortranspec.c b/gcc/fortran/gfortranspec.c index 33e6e572b54a..a8e87d8e5c3c 100644 --- a/gcc/fortran/gfortranspec.c +++ b/gcc/fortran/gfortranspec.c @@ -441,7 +441,7 @@ int lang_specific_pre_link (void) { if (library) - do_spec ("%:include(libgfortran.spec)"); + do_spec ("%:include(libgfortran.spec)", 0); return 0; } diff --git a/gcc/gcc.c b/gcc/gcc.c index 4f57765b012a..940e87fb2064 100644 --- a/gcc/gcc.c +++ b/gcc/gcc.c @@ -409,6 +409,7 @@ static const char *replace_extension_spec_func (int, const char **); static const char *greater_than_spec_func (int, const char **); static const char *debug_level_greater_than_spec_func (int, const char **); static const char *find_fortran_preinclude_file (int, const char **); +static const char *add_omp_infile_spec_func (int, const char **); static char *convert_white_space (char *); /* The Specs Language @@ -1222,6 +1223,9 @@ static const char *const multilib_defaults_raw[] = MULTILIB_DEFAULTS; static const char *const driver_self_specs[] = { "%{fdump-final-insns:-fdump-final-insns=.} %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) @@ -3747,7 +3763,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 @@ -3774,7 +3791,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; @@ -4033,7 +4051,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); set_source_date_epoch_envvar (); return true; @@ -4094,17 +4113,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; @@ -4121,19 +4140,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: @@ -4260,7 +4281,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 @@ -4294,9 +4316,12 @@ 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; } @@ -4596,7 +4621,7 @@ process_command (unsigned int decoded_options_count, error ("%s: %m", fname + resp); } else - add_infile (arg, spec_lang); + add_infile (arg, spec_lang, 0); free (fname); continue; @@ -4746,7 +4771,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; } @@ -4757,7 +4783,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); } /* Decide if undefined variable references are allowed in specs. */ @@ -4978,13 +5004,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, NULL); /* Force out any unfinished command. @@ -5144,7 +5172,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: @@ -6723,6 +6752,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; @@ -7829,9 +7866,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. */ @@ -8145,7 +8179,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; @@ -8160,7 +8195,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]; @@ -8315,7 +8351,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); @@ -9992,6 +10028,53 @@ find_fortran_preinclude_file (int argc, const char **argv) return result; } +/* 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); + + /* Nothing to do if we're not actually offloading. */ + if (!ENABLE_OFFLOADING) + return NULL; + 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_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); + + 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 a/gcc/gcc.h b/gcc/gcc.h index a0a1d94c6e64..21bd036341bf 100644 --- a/gcc/gcc.h +++ b/gcc/gcc.h @@ -69,7 +69,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 set_input (const char *); diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp new file mode 100644 index 000000000000..082ce68c654a --- /dev/null +++ b/libgomp/ChangeLog.omp @@ -0,0 +1,63 @@ +2018-05-20 Thomas Schwinge + + 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-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. + +2017-05-14 Thomas Schwinge + + PR libgomp/81886 + * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Adapt. + +2015-08-20 Thomas Schwinge + Joseph Myers + + PR libgomp/81886 + * plugin/configfrag.ac (tgt_name): Do not set. + (offload_targets): Separate with colons not commas. + * config.h.in, configure: Regenerate. + * libgomp.h (gomp_offload_target_available_p): New prototype. + * libgomp.map (GOACC_2.0.GOMP_4_BRANCH): Add + GOMP_set_offload_targets. + * libgomp_g.h (GOMP_set_offload_targets): New prototype. + * oacc-init.c (resolve_device): Use + gomp_offload_target_available_p. + * target.c (resolve_device): Use host fallback when offload data + not available. + (gomp_offload_target_available_p, offload_target_to_plugin_name) + (gomp_offload_targets, gomp_offload_targets_init) + (GOMP_set_offload_targets, gomp_plugin_prefix) + (gomp_plugin_suffix): New. + (gomp_load_plugin_for_device): Add gomp_debug call. + (gomp_target_init): Use gomp_offload_targets instead of + OFFLOAD_TARGETS. Handle and rewrie colon-separated string. + * testsuite/lib/libgomp.exp: Expect offload targets to be + colon-separated. Adjust matching of offload targets. + (libgomp_init) + (check_effective_target_openacc_nvidia_accel_configured) + (check_effective_target_openacc_host_selected): Adjust checks of + offload target names. + * testsuite/libgomp.oacc-c++/c++.exp: Adjust set of offload + targets. Use -foffload instead of setenv ACC_DEVICE_TYPE. + * testsuite/libgomp.oacc-c/c.exp: Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. diff --git a/libgomp/config.h.in b/libgomp/config.h.in index 73f1b12805e3..449cc8506878 100644 --- a/libgomp/config.h.in +++ b/libgomp/config.h.in @@ -146,8 +146,8 @@ */ #undef LT_OBJDIR -/* Define to offload plugins, separated by commas. */ -#undef OFFLOAD_PLUGINS +/* Define to offload targets, separated by colons. */ +#undef OFFLOAD_TARGETS /* Name of package */ #undef PACKAGE diff --git a/libgomp/configure b/libgomp/configure index b4bc4f436289..529a2b4711d5 100755 --- a/libgomp/configure +++ b/libgomp/configure @@ -15743,8 +15743,8 @@ rm -f core conftest.err conftest.$ac_objext \ offload_plugins=$tgt_plugin offload_targets=$tgt else - offload_plugins=$offload_plugins,$tgt_plugin - offload_targets=$offload_targets,$tgt + offload_plugins=$offload_plugins:$tgt_plugin + offload_targets=$offload_targets:$tgt fi # Configure additional search paths. if test "$tgt_plugin" = hsa; then @@ -15761,7 +15761,7 @@ rm -f core conftest.err conftest.$ac_objext \ fi cat >>confdefs.h <<_ACEOF -#define OFFLOAD_PLUGINS "$offload_plugins" +#define OFFLOAD_TARGETS "$offload_targets" _ACEOF if test $PLUGIN_NVPTX = 1; then diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index afea659445d8..2483817eb9ad 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1067,6 +1067,7 @@ extern void gomp_init_device (struct gomp_device_descr *); extern void gomp_free_memmap (struct splay_tree_s *); extern void gomp_unload_device (struct gomp_device_descr *); extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key); +extern bool gomp_offload_target_available_p (int); /* work.c */ diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index d8e2fd1818b9..b633df438529 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -494,6 +494,11 @@ GOACC_2.0.1 { GOACC_parallel_keyed; } GOACC_2.0; +GOACC_2.0.GOMP_4_BRANCH { + global: + GOMP_set_offload_targets; +} GOACC_2.0.1; + GOMP_PLUGIN_1.0 { global: GOMP_PLUGIN_malloc; diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index 32a9d8aade90..5c4bd8d34971 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -334,6 +334,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_ext (int, void (*) (void *), size_t, void **, size_t *, diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c index f30cf2f81d8d..ac1464773551 100644 --- a/libgomp/oacc-init.c +++ b/libgomp/oacc-init.c @@ -87,6 +87,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 @@ -103,6 +105,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); } } @@ -114,6 +118,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) @@ -122,7 +128,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 +156,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 +178,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 +192,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) { @@ -190,6 +202,7 @@ resolve_device (acc_device_t d, bool fail_is_error) gomp_fatal ("device type %s not supported", name_of_acc_device_t (d)); } + gomp_debug (0, " %s: %d: %p\n", __FUNCTION__, (int) d, dispatchers[d]); return dispatchers[d]; } diff --git a/libgomp/openacc.h b/libgomp/openacc.h index 1bbe6c90e7fc..fa5b3ae5dbcc 100644 --- a/libgomp/openacc.h +++ b/libgomp/openacc.h @@ -55,6 +55,8 @@ typedef enum acc_device_t { /* acc_device_host_nonshm = 3 removed. */ acc_device_not_host = 4, acc_device_nvidia = 5, + /* not supported */ _acc_device_intel_mic = 6, + /* not supported */ _acc_device_hsa = 7, _ACC_device_hwm, /* Ensure enumeration is layout compatible with int. */ _ACC_highest = __INT_MAX__, diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac index 9718ac752e2d..13ca26f47d99 100644 --- a/libgomp/plugin/configfrag.ac +++ b/libgomp/plugin/configfrag.ac @@ -249,8 +249,8 @@ if test x"$enable_offload_targets" != x; then offload_plugins=$tgt_plugin offload_targets=$tgt else - offload_plugins=$offload_plugins,$tgt_plugin - offload_targets=$offload_targets,$tgt + offload_plugins=$offload_plugins:$tgt_plugin + offload_targets=$offload_targets:$tgt fi # Configure additional search paths. if test "$tgt_plugin" = hsa; then @@ -265,8 +265,8 @@ if test x"$enable_offload_targets" != x; then fi done fi -AC_DEFINE_UNQUOTED(OFFLOAD_PLUGINS, "$offload_plugins", - [Define to offload plugins, separated by commas.]) +AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets", + [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 a/libgomp/target.c b/libgomp/target.c index 31148003d0a5..4bfebf481c10 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -106,13 +106,18 @@ 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) + gomp_debug (0, "%s (%d)\n", __FUNCTION__, device); + + 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; @@ -127,6 +132,13 @@ resolve_device (int device_id) } gomp_mutex_unlock (&devices[device_id].lock); + /* 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; + + gomp_debug (0, " %s (%d): %d\n", __FUNCTION__, device, device_id); return &devices[device_id]; } @@ -1426,6 +1438,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)) { @@ -1481,6 +1496,52 @@ gomp_free_memmap (struct splay_tree_s *mem_map) } } +/* 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) +{ + gomp_debug (0, "%s (%d)\n", __FUNCTION__, 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->state == GOMP_DEVICE_INITIALIZED) + available = true; + gomp_mutex_unlock (&devicep->lock); + } + + if (!available) + { + 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); + } + + gomp_debug (0, " %s (%d): %d\n", __FUNCTION__, type, (int) available); + return available; +} + /* Host fallback for GOMP_target{,_ext} routines. */ static void @@ -2588,6 +2649,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); @@ -2710,6 +2773,190 @@ gomp_target_fini (void) } } +/* Helper, to translate from an offload target to the corresponding plugin name. */ +/* TODO: this duplicates the logic/information that we already have in + 'offload_targets' vs. 'offload_plugins' variables, + 'libgomp/plugin/configfrag.ac'. */ + +static const char * +offload_target_to_plugin_name (const char *offload_target) +{ + if (strstr (offload_target, "-intelmic") != NULL) + return "intelmic"; + else if (strncmp (offload_target, "nvptx", 5) == 0) + return "nvptx"; + else if (strncmp (offload_target, "hsa", 3) == 0) + return "hsa"; + else + gomp_fatal ("Unknown offload target: %s", offload_target); +} + +/* 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_set = false; +static bool gomp_offload_targets_malloced = false; + +/* 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\"): %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 for offloading. It parses the list of offload plugins, and tries to load these. On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP @@ -2717,11 +2964,12 @@ gomp_target_fini (void) 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; @@ -2729,52 +2977,60 @@ gomp_target_init (void) num_devices = 0; devices = NULL; - cur = OFFLOAD_PLUGINS; + cur = gomp_offload_targets; if (*cur) do { - struct gomp_device_descr current_device; - size_t prefix_len, suffix_len, cur_len; - - next = strchr (cur, ','); - - prefix_len = strlen (prefix); - cur_len = next ? next - cur : strlen (cur); - suffix_len = strlen (suffix); - - plugin_name = (char *) malloc (prefix_len + cur_len + suffix_len + 1); - if (!plugin_name) - { - num_devices = 0; - break; - } - - memcpy (plugin_name, prefix, prefix_len); - memcpy (plugin_name + prefix_len, cur, cur_len); - memcpy (plugin_name + prefix_len + cur_len, suffix, suffix_len + 1); + 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.state = GOMP_DEVICE_UNINITIALIZED; 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; @@ -2788,18 +3044,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 a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 14d9b5f13054..a0fe4071cf9a 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -120,7 +120,7 @@ proc libgomp_init { args } { # Add liboffloadmic build directory in LD_LIBRARY_PATH to support # Intel MIC offloading testing. global offload_plugins - if { [string match "*,intelmic,*" ",$offload_plugins,"] } { + if { [string match "*:intelmic:*" ":$offload_plugins:"] } { append always_ld_library_path ":${blddir}/../liboffloadmic/.libs" append always_ld_library_path ":${blddir}/../liboffloadmic/plugin/.libs" # libstdc++ is required by liboffloadmic @@ -373,22 +373,23 @@ proc check_effective_target_openacc_nvidia_accel_present { } { } "" ] } -# Return 1 if at least one Nvidia GPU is accessible, and the OpenACC 'nvidia' -# device type is selected. +# Return 1 if at least one Nvidia GPU is accessible, and 'nvptx' offloading is +# selected by 'global offload_target'. proc check_effective_target_openacc_nvidia_accel_selected { } { if { ![check_effective_target_openacc_nvidia_accel_present] } { return 0; } - global openacc_device_type - return [string match "nvidia" $openacc_device_type] + global offload_target + return [string match "nvptx*" $offload_target] } -# Return 1 if the OpenACC 'host' device type is selected. +# Return 1 if explicit host-fallback execution is selected by 'global +# offload_target'. proc check_effective_target_openacc_host_selected { } { - global openacc_device_type - return [string match "host" $openacc_device_type] + global offload_target + return [string match "disable" $offload_target] } # Return 1 if the selected OMP device is actually a HSA device diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp index dcefa792ca40..b8b44518b861 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -78,11 +78,13 @@ if { $lang_test_file_found } { } # Test with all available offload targets, and with offloading disabled. - foreach offload_target [concat [split $offload_targets ","] "disable"] { - global openacc_device_type - set openacc_device_type [offload_target_to_openacc_device_type $offload_target] - set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" + set SAVE_ALWAYS_CFLAGS "$ALWAYS_CFLAGS" + global offload_target + foreach offload_target [concat [split $offload_targets ":"] "disable"] { + set ALWAYS_CFLAGS "$SAVE_ALWAYS_CFLAGS" + set openacc_device_type [offload_target_to_openacc_device_type $offload_target] + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=\"$offload_target\"" switch $openacc_device_type { "" { unsupported "$subdir $offload_target offloading" @@ -90,6 +92,9 @@ if { $lang_test_file_found } { } host { set acc_mem_shared 1 + + # Special case: pass the empty string instead of "disable". + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=\"\"" } nvidia { if { ![check_effective_target_openacc_nvidia_accel_present] } { @@ -116,8 +121,6 @@ if { $lang_test_file_found } { # handling in test cases, by default only build for the offload target # that we're actually going to test. set tagopt "$tagopt -foffload=$offload_target" - # Force usage of the corresponding OpenACC device type. - setenv ACC_DEVICE_TYPE $openacc_device_type # To get better test coverage for device-specific code that is only # ever used in offloading configurations, we'd like more thorough diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c index bfcb67d58b60..758b1fcb366c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c @@ -14,7 +14,7 @@ int main () int expect = 1; -#if ACC_DEVICE_TYPE_host +#ifdef ACC_DEVICE_TYPE_host expect = 0; #endif diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c index 8112745bcb86..0270d06fcdaa 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c @@ -37,7 +37,7 @@ main (int argc, char *argv[]) } -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host /* Offloaded. */ @@ -49,7 +49,7 @@ main (int argc, char *argv[]) abort (); if (!acc_on_device (acc_device_not_host)) abort (); -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia if (!acc_on_device (acc_device_nvidia)) abort (); #else diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h index 9db236c83629..ca29f0889e4e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h @@ -1,5 +1,5 @@ -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia #pragma acc routine nohost static int clock (void) diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp index 55cd40f1e991..aba6665b5cb0 100644 --- a/libgomp/testsuite/libgomp.oacc-c/c.exp +++ b/libgomp/testsuite/libgomp.oacc-c/c.exp @@ -40,12 +40,12 @@ set_ld_library_path_env_vars # Test with all available offload targets, and with offloading disabled. set SAVE_ALWAYS_CFLAGS "$ALWAYS_CFLAGS" -foreach offload_target [concat [split $offload_targets ","] "disable"] { +global offload_target +foreach offload_target [concat [split $offload_targets ":"] "disable"] { set ALWAYS_CFLAGS "$SAVE_ALWAYS_CFLAGS" - global openacc_device_type - set openacc_device_type [offload_target_to_openacc_device_type $offload_target] - set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" + set openacc_device_type [offload_target_to_openacc_device_type $offload_target] + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=\"$offload_target\"" switch $openacc_device_type { "" { unsupported "$subdir $offload_target offloading" @@ -53,6 +53,9 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { } host { set acc_mem_shared 1 + + # Special case: pass the empty string instead of "disable". + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=\"\"" } nvidia { if { ![check_effective_target_openacc_nvidia_accel_present] } { @@ -79,8 +82,6 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { # handling in test cases, by default only build for the offload target # that we're actually going to test. set tagopt "$tagopt -foffload=$offload_target" - # Force usage of the corresponding OpenACC device type. - setenv ACC_DEVICE_TYPE $openacc_device_type # To get better test coverage for device-specific code that is only # ever used in offloading configurations, we'd like more thorough diff --git a/libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c b/libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c new file mode 100644 index 000000000000..b62a587ec08f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c @@ -0,0 +1,119 @@ +/* Test what happens for repeated GOMP_set_offload_targets calls, which happens + when shared libraries are involved, for example. As in the libgomp + testsuite infrastructure, it is difficult to build and link against shared + libraries, we simulate that by replicating some relevant + GOMP_set_offload_targets calls. */ + +#include +#include +#include +#include +#include "libgomp_g.h" + +int main () +{ + /* Before getting here, GOMP_set_offload_targets already got called via a + constructor. */ + + bool acc_device_types_requested[_ACC_device_hwm]; + for (int i = 0; i < _ACC_device_hwm; ++i) + acc_device_types_requested[i] = false; + + /* We're building for only one offload target ("-foffload=[...]") which is + the following. */ + const char *offload_target_requested; + acc_device_t acc_device_type_requested; +#if defined ACC_DEVICE_TYPE_nvidia + offload_target_requested = ACC_DEVICE_TYPE_nvidia; + acc_device_type_requested = acc_device_nvidia; +#elif defined ACC_DEVICE_TYPE_host + offload_target_requested = ACC_DEVICE_TYPE_host; + acc_device_type_requested = acc_device_host; +#else +# error Not ported to this ACC_DEVICE_TYPE +#endif + acc_device_types_requested[acc_device_type_requested] = true; + +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + /* Call again; will have no noticeable difference. */ + GOMP_set_offload_targets (offload_target_requested); +#endif + +#ifdef OFFLOAD_TARGETS_ADD_EARLY + /* Request a (non-existing) offloading target (which will result in a + non-fatal diagnostic). */ + GOMP_set_offload_targets (OFFLOAD_TARGETS_ADD); +#endif + +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + /* Call again; will have no noticeable difference. */ + GOMP_set_offload_targets (offload_target_requested); + char *s; + { + size_t len = 3 * (strlen (offload_target_requested) + 1); +# ifdef OFFLOAD_TARGETS_ADD_EARLY + len += 3 * (strlen (OFFLOAD_TARGETS_ADD) + 1); +# endif + s = malloc (len); + if (s == NULL) + __builtin_abort (); + size_t len_; +# ifndef OFFLOAD_TARGETS_ADD_EARLY + len_ = sprintf (s, "%s:%s:%s", + offload_target_requested, + offload_target_requested, + offload_target_requested); +# else + len_ = sprintf (s, "%s:%s:%s:%s:%s:%s", + offload_target_requested, + offload_target_requested, + OFFLOAD_TARGETS_ADD, + OFFLOAD_TARGETS_ADD, + offload_target_requested, + OFFLOAD_TARGETS_ADD); +# endif + if (len_ + 1 != len) + __builtin_abort (); + GOMP_set_offload_targets (s); + } +#endif + + /* Calling acc_get_num_devices will implicitly initialize offloading. */ +#if defined OFFLOAD_TARGETS_ADD_EARLY + fprintf (stderr, "CheCKpOInT1\n"); +#endif + /* acc_device_host is always available. */ + if ((acc_get_num_devices (acc_device_host) > 0) == false) + __builtin_abort (); +#if defined OFFLOAD_TARGETS_ADD_EARLY + fprintf (stderr, "WrONg WAy1\n"); +#endif + for (acc_device_t acc_device_type = acc_device_not_host + 1; + acc_device_type < _ACC_device_hwm; + ++acc_device_type) + { + /* The requested device type must be available. Any other device types + must not be available. */ + if ((acc_get_num_devices (acc_device_type) > 0) + != acc_device_types_requested[acc_device_type]) + __builtin_abort (); + } + +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + /* Request the same again; will have no noticeable difference. */ + GOMP_set_offload_targets (offload_target_requested); +#endif +#if defined OFFLOAD_TARGETS_ADD_LATE + fprintf (stderr, "CheCKpOInT2\n"); + GOMP_set_offload_targets (OFFLOAD_TARGETS_ADD); + fprintf (stderr, "WrONg WAy2\n"); +#endif +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + GOMP_set_offload_targets (s); + + /* Implementation defail: OK to "free (s)", in this case. */ + free (s); +#endif + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c b/libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c new file mode 100644 index 000000000000..977c559314e1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c @@ -0,0 +1,2 @@ +#define OFFLOAD_TARGETS_SAME_AGAIN +#include "offload-targets-1.c" diff --git a/libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c b/libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c new file mode 100644 index 000000000000..1eb080be33fa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c @@ -0,0 +1,10 @@ +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_EARLY +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT1(\n|\r\n|\r)+" } + { dg-output "libgomp: Unknown offload target: XYZ(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git a/libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c b/libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c new file mode 100644 index 000000000000..2bb7204ed419 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c @@ -0,0 +1,11 @@ +#define OFFLOAD_TARGETS_SAME_AGAIN +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_EARLY +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT1(\n|\r\n|\r)+" } + { dg-output "libgomp: Unknown offload target: XYZ(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git a/libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c b/libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c new file mode 100644 index 000000000000..8ba0792af49b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c @@ -0,0 +1,10 @@ +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_LATE +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT2(\n|\r\n|\r)+" } + { dg-output "libgomp: Can't satisfy request for offload targets: XYZ; have loaded: \[a-z-\]*(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git a/libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c b/libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c new file mode 100644 index 000000000000..4b15582106a1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c @@ -0,0 +1,11 @@ +#define OFFLOAD_TARGETS_SAME_AGAIN +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_LATE +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT2(\n|\r\n|\r)+" } + { dg-output "libgomp: Can't satisfy request for offload targets: XYZ; have loaded: \[a-z-\]*(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 index d6c67a0c31aa..6a82385e6e6c 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 @@ -25,7 +25,7 @@ if (acc_on_device (acc_device_nvidia)) STOP 8 !$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)) STOP 8 if (acc_on_device (acc_device_none)) STOP 9 if (acc_on_device (acc_device_host)) STOP 10 if (.not. acc_on_device (acc_device_not_host)) STOP 11 -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia if (.not. acc_on_device (acc_device_nvidia)) STOP 12 #else if (acc_on_device (acc_device_nvidia)) STOP 13 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f index 75e24509ce9b..1b9f9ac2a1ab 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f @@ -25,7 +25,7 @@ !$ACC END PARALLEL -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host ! Offloaded. @@ -33,7 +33,7 @@ IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) STOP 9 IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) STOP 10 IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) STOP 11 -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) STOP 12 #else IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) STOP 13 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f index 908d185f40c7..82bf95459413 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f @@ -25,7 +25,7 @@ !$ACC END PARALLEL -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host ! Offloaded. @@ -33,7 +33,7 @@ IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) STOP 9 IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) STOP 10 IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) STOP 11 -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) STOP 12 #else IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) STOP 13 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp index af25a22a522b..fb54e486c87b 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp +++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp @@ -66,11 +66,10 @@ if { $lang_test_file_found } { set_ld_library_path_env_vars # Test with all available offload targets, and with offloading disabled. - foreach offload_target [concat [split $offload_targets ","] "disable"] { - global openacc_device_type + global offload_target + foreach offload_target [concat [split $offload_targets ":"] "disable"] { set openacc_device_type [offload_target_to_openacc_device_type $offload_target] - set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" - + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=\"$offload_target\"" switch $openacc_device_type { "" { unsupported "$subdir $offload_target offloading" @@ -78,6 +77,9 @@ if { $lang_test_file_found } { } host { set acc_mem_shared 1 + + # Special case: pass the empty string instead of "disable". + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=\"\"" } nvidia { if { ![check_effective_target_openacc_nvidia_accel_present] } { @@ -98,8 +100,6 @@ if { $lang_test_file_found } { # handling in test cases, by default only build for the offload target # that we're actually going to test. set tagopt "$tagopt -foffload=$offload_target" - # Force usage of the corresponding OpenACC device type. - setenv ACC_DEVICE_TYPE $openacc_device_type # 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