@@ -1026,12 +1026,20 @@ AC_SUBST(real_target_noncanonical)
AC_SUBST(accel_dir_suffix)
for tgt in `echo $enable_offload_targets | sed 's/,/ /g'`; do
+ tgt_dir=`echo $tgt | sed -n 's/.*=//p'`
tgt=`echo $tgt | sed 's/=.*//'`
if echo "$tgt" | grep "^hsa" > /dev/null ; then
enable_hsa=1
else
enable_offloading=1
+ if test -n "$tgt_dir"; then
+ omp_device_property="${tgt_dir}/lib/gcc/\$(real_target_noncanonical)/\$(version)/accel/${tgt}/omp-device-properties"
+ else
+ omp_device_property="\$(libsubdir)/accel/${tgt}/omp-device-properties"
+ fi
+ omp_device_properties="${omp_device_properties} ${tgt}=${omp_device_property}"
+ omp_device_property_deps="${omp_device_property_deps} ${omp_device_property}"
fi
if test x"$offload_targets" = x; then
@@ -1040,6 +1048,9 @@ for tgt in `echo $enable_offload_targets
offload_targets="$offload_targets,$tgt"
fi
done
+AC_SUBST(omp_device_properties)
+AC_SUBST(omp_device_property_deps)
+
AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets",
[Define to offload targets, separated by commas.])
if test x"$enable_offloading" != x; then
@@ -2645,7 +2645,7 @@ generated_files = config.h tm.h $(TM_P_H
common/common-target-hooks-def.h pass-instances.def \
c-family/c-target-hooks-def.h d/d-target-hooks-def.h \
params.list params.options case-cfn-macros.h \
- cfn-operators.pd
+ cfn-operators.pd omp-device-properties.h
#
# How to compile object files to run on the build machine.
@@ -2854,6 +2854,30 @@ $(genprog:%=build/gen%$(build_exeext)):
+$(LINKER_FOR_BUILD) $(BUILD_LINKERFLAGS) $(BUILD_LDFLAGS) -o $@ \
$(filter-out $(BUILD_LIBDEPS), $^) $(BUILD_LIBS)
+omp-general.o: omp-device-properties.h
+
+omp_device_properties = @omp_device_properties@
+omp-device-properties.h: s-omp-device-properties-h ; @true
+s-omp-device-properties-h: @omp_device_property_deps@
+ -rm -f tmp-omp-device-properties.h; \
+ for kind in kind arch isa; do \
+ echo 'const char omp_offload_device_'$${kind}'[] = ' \
+ >> tmp-omp-device-properties.h; \
+ for prop in none $(omp_device_properties); do \
+ [ "$$prop" = "none" ] && continue; \
+ tgt=`echo "$$prop" | sed 's/=.*$$//'`; \
+ props=`echo "$$prop" | sed 's/.*=//'`; \
+ echo "\"$$tgt\\0\"" >> tmp-omp-device-properties.h; \
+ sed -n 's/^'$${kind}': //p' $${props} \
+ | sed 's/[[:blank:]]/ /g;s/ */ /g;s/^ //;s/ $$//;s/ /\\0/g;s/^/"/;s/$$/\\0\\0"/' \
+ >> tmp-omp-device-properties.h; \
+ done; \
+ echo '"";' >> tmp-omp-device-properties.h; \
+ done; \
+ $(SHELL) $(srcdir)/../move-if-change tmp-omp-device-properties.h \
+ omp-device-properties.h
+ $(STAMP) s-omp-device-properties-h
+
# Generated source files for gengtype. Prepend inclusion of
# config.h/bconfig.h because AIX requires _LARGE_FILES to be defined before
# any system header is included.
@@ -3452,6 +3476,10 @@ ifeq ($(enable_plugin),yes)
install: install-plugin
endif
+ifeq ($(enable_as_accelerator),yes)
+install: install-omp-device-properties
+endif
+
install-strip: override INSTALL_PROGRAM = $(INSTALL_STRIP_PROGRAM)
ifneq ($(STRIP),)
install-strip: STRIPPROG = $(STRIP)
@@ -3637,6 +3665,11 @@ install-driver: installdirs xgcc$(exeext
fi; \
fi
+# Install omp-device-properties file for accelerator compilers.
+install-omp-device-properties: omp-device-properties installdirs
+ $(INSTALL_DATA) omp-device-properties \
+ $(DESTDIR)$(libsubdir)/omp-device-properties
+
# Install the info files.
# $(INSTALL_DATA) might be a relative pathname, so we can't cd into srcdir
# to do the install.
@@ -1669,6 +1669,21 @@ int, (void), NULL)
HOOK_VECTOR_END (simt)
+/* Functions relating to OpenMP. */
+#undef HOOK_PREFIX
+#define HOOK_PREFIX "TARGET_OMP_"
+HOOK_VECTOR (TARGET_OMP, omp)
+
+DEFHOOK
+(device_kind_arch_isa,
+"Return 1 if @var{trait} @var{name} is present in the OpenMP context's\n\
+device trait set, return 0 if not present in any OpenMP context in the\n\
+whole translation unit, or -1 if not present in the current OpenMP context\n\
+but might be present in another OpenMP context in the same TU.",
+int, (enum omp_device_kind_arch_isa trait, const char *name), NULL)
+
+HOOK_VECTOR_END (omp)
+
/* Functions relating to openacc. */
#undef HOOK_PREFIX
#define HOOK_PREFIX "TARGET_GOACC_"
@@ -211,6 +211,13 @@ typedef vec<poly_uint64> vector_sizes;
automatically freed. */
typedef auto_vec<poly_uint64, 8> auto_vector_sizes;
+/* First argument of targetm.omp.device_kind_arch_isa. */
+enum omp_device_kind_arch_isa {
+ omp_device_kind,
+ omp_device_arch,
+ omp_device_isa
+};
+
/* The target structure. This holds all the backend hooks. */
#define DEFHOOKPOD(NAME, DOC, TYPE, INIT) TYPE NAME;
#define DEFHOOK(NAME, DOC, TYPE, PARAMS, INIT) TYPE (* NAME) PARAMS;
@@ -4205,6 +4205,8 @@ address; but often a machine-dependent
@hook TARGET_SIMT_VF
+@hook TARGET_OMP_DEVICE_KIND_ARCH_ISA
+
@hook TARGET_GOACC_VALIDATE_DIMS
@hook TARGET_GOACC_DIM_LIMIT
@@ -40,6 +40,7 @@ along with GCC; see the file COPYING3.
#include "symbol-summary.h"
#include "hsa-common.h"
#include "tree-pass.h"
+#include "omp-device-properties.h"
enum omp_requires omp_requires_mask;
@@ -537,7 +538,7 @@ omp_max_simt_vf (void)
{
if (!strncmp (c, "nvptx", strlen ("nvptx")))
return 32;
- else if ((c = strchr (c, ',')))
+ else if ((c = strchr (c, ':')))
c++;
}
return 0;
@@ -571,6 +572,79 @@ omp_constructor_traits_to_codes (tree ct
return nconstructs;
}
+/* Return true if PROP is possibly present in one of the offloading target's
+ OpenMP contexts. The format of PROPS string is always offloading target's
+ name terminated by '\0', followed by properties for that offloading
+ target separated by '\0' and terminated by another '\0'. The strings
+ are created from omp-device-properties installed files of all configured
+ offloading targets. */
+
+static bool
+omp_offload_device_kind_arch_isa (const char *props, const char *prop)
+{
+ const char *names = getenv ("OFFLOAD_TARGET_NAMES");
+ if (names == NULL || *names == '\0')
+ return false;
+ while (*props != '\0')
+ {
+ size_t name_len = strlen (props);
+ bool matches = false;
+ for (const char *c = names; c; )
+ {
+ if (strncmp (props, c, name_len) == 0
+ && (c[name_len] == '\0'
+ || c[name_len] == ':'
+ || c[name_len] == '='))
+ {
+ matches = true;
+ break;
+ }
+ else if ((c = strchr (c, ':')))
+ c++;
+ }
+ props = props + name_len + 1;
+ while (*props != '\0')
+ {
+ if (matches && strcmp (props, prop) == 0)
+ return true;
+ props = strchr (props, '\0') + 1;
+ }
+ props++;
+ }
+ return false;
+}
+
+/* Return true if the current code location is or might be offloaded.
+ Return true in declare target functions, or when nested in a target
+ region or when unsure, return false otherwise. */
+
+static bool
+omp_maybe_offloaded (void)
+{
+ if (!hsa_gen_requested_p ())
+ {
+ if (!ENABLE_OFFLOADING)
+ return false;
+ const char *names = getenv ("OFFLOAD_TARGET_NAMES");
+ if (names == NULL || *names == '\0')
+ return false;
+ }
+ if (symtab->state == PARSING)
+ /* Maybe. */
+ return true;
+ if (current_function_decl
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (current_function_decl)))
+ return true;
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
+ {
+ enum tree_code construct = OMP_TARGET;
+ if (omp_construct_selector_matches (&construct, 1))
+ return true;
+ }
+ return false;
+}
+
/* Return 1 if context selector matches the current OpenMP context, 0
if it does not and -1 if it is unknown and need to be determined later.
Some properties can be checked right away during parsing (this routine),
@@ -667,8 +741,45 @@ omp_context_selector_matches (tree ctx)
return 0;
}
if (set == 'd' && !strcmp (sel, "arch"))
- /* For now, need a target hook. */
- ret = -1;
+ for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
+ {
+ const char *arch = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
+ int r = 0;
+ if (targetm.omp.device_kind_arch_isa != NULL)
+ r = targetm.omp.device_kind_arch_isa (omp_device_arch,
+ arch);
+ if (r == 0 || (r == -1 && symtab->state != PARSING))
+ {
+ /* If we are or might be in a target region or
+ declare target function, need to take into account
+ also offloading values. */
+ if (!omp_maybe_offloaded ())
+ return 0;
+ if (strcmp (arch, "hsa") == 0
+ && hsa_gen_requested_p ())
+ {
+ ret = -1;
+ continue;
+ }
+ if (ENABLE_OFFLOADING)
+ {
+ const char *arches = omp_offload_device_arch;
+ if (omp_offload_device_kind_arch_isa (arches,
+ arch))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ return 0;
+ }
+ else if (r == -1)
+ ret = -1;
+ /* If arch matches on the host, it still might not match
+ in the offloading region. */
+ else if (omp_maybe_offloaded ())
+ ret = -1;
+ }
break;
case 'u':
if (set == 'i' && !strcmp (sel, "unified_address"))
@@ -729,57 +840,92 @@ omp_context_selector_matches (tree ctx)
const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
if (!strcmp (prop, "any"))
continue;
- if (!strcmp (prop, "fpga"))
- return 0; /* Right now GCC doesn't support any fpgas. */
if (!strcmp (prop, "host"))
{
- if (ENABLE_OFFLOADING || hsa_gen_requested_p ())
+ if (omp_maybe_offloaded ())
ret = -1;
continue;
}
if (!strcmp (prop, "nohost"))
{
- if (ENABLE_OFFLOADING || hsa_gen_requested_p ())
+ if (omp_maybe_offloaded ())
ret = -1;
else
return 0;
continue;
}
- if (!strcmp (prop, "cpu") || !strcmp (prop, "gpu"))
+ int r = 0;
+ if (targetm.omp.device_kind_arch_isa != NULL)
+ r = targetm.omp.device_kind_arch_isa (omp_device_kind,
+ prop);
+ else
+ r = strcmp (prop, "cpu") == 0;
+ if (r == 0 || (r == -1 && symtab->state != PARSING))
{
- bool maybe_gpu = false;
- if (hsa_gen_requested_p ())
- maybe_gpu = true;
- else if (ENABLE_OFFLOADING)
- for (const char *c = getenv ("OFFLOAD_TARGET_NAMES");
- c; )
- {
- if (!strncmp (c, "nvptx", strlen ("nvptx"))
- || !strncmp (c, "amdgcn", strlen ("amdgcn")))
- {
- maybe_gpu = true;
- break;
- }
- else if ((c = strchr (c, ',')))
- c++;
- }
- if (!maybe_gpu)
+ /* If we are or might be in a target region or
+ declare target function, need to take into account
+ also offloading values. */
+ if (!omp_maybe_offloaded ())
+ return 0;
+ if (strcmp (prop, "gpu") == 0
+ && hsa_gen_requested_p ())
{
- if (prop[0] == 'g')
- return 0;
+ ret = -1;
+ continue;
}
- else
- ret = -1;
- continue;
+ if (ENABLE_OFFLOADING)
+ {
+ const char *kinds = omp_offload_device_kind;
+ if (omp_offload_device_kind_arch_isa (kinds, prop))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ return 0;
}
- /* Any other kind doesn't match. */
- return 0;
+ else if (r == -1)
+ ret = -1;
+ /* If kind matches on the host, it still might not match
+ in the offloading region. */
+ else if (omp_maybe_offloaded ())
+ ret = -1;
}
break;
case 'i':
if (set == 'd' && !strcmp (sel, "isa"))
- /* For now, need a target hook. */
- ret = -1;
+ for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
+ {
+ const char *isa = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
+ int r = 0;
+ if (targetm.omp.device_kind_arch_isa != NULL)
+ r = targetm.omp.device_kind_arch_isa (omp_device_isa,
+ isa);
+ if (r == 0 || (r == -1 && symtab->state != PARSING))
+ {
+ /* If we are or might be in a target region or
+ declare target function, need to take into account
+ also offloading values. */
+ if (!omp_maybe_offloaded ())
+ return 0;
+ if (ENABLE_OFFLOADING)
+ {
+ const char *isas = omp_offload_device_isa;
+ if (omp_offload_device_kind_arch_isa (isas, isa))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ return 0;
+ }
+ else if (r == -1)
+ ret = -1;
+ /* If isa matches on the host, it still might not match
+ in the offloading region. */
+ else if (omp_maybe_offloaded ())
+ ret = -1;
+ }
break;
case 'c':
if (set == 'u' && !strcmp (sel, "condition"))
@@ -19,6 +19,9 @@ along with GCC; see the file COPYING3.
#ifndef GCC_I386_OPTIONS_H
#define GCC_I386_OPTIONS_H
+extern int ix86_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
+ const char *name);
+
char *ix86_target_string (HOST_WIDE_INT isa, HOST_WIDE_INT isa2,
int flags, int flags2,
const char *arch, const char *tune,
@@ -23035,12 +23035,13 @@ ix86_run_selftests (void)
ix86_simd_clone_compute_vecsize_and_simdlen
#undef TARGET_SIMD_CLONE_ADJUST
-#define TARGET_SIMD_CLONE_ADJUST \
- ix86_simd_clone_adjust
+#define TARGET_SIMD_CLONE_ADJUST ix86_simd_clone_adjust
#undef TARGET_SIMD_CLONE_USABLE
-#define TARGET_SIMD_CLONE_USABLE \
- ix86_simd_clone_usable
+#define TARGET_SIMD_CLONE_USABLE ix86_simd_clone_usable
+
+#undef TARGET_OMP_DEVICE_KIND_ARCH_ISA
+#define TARGET_OMP_DEVICE_KIND_ARCH_ISA ix86_omp_device_kind_arch_isa
#undef TARGET_FLOAT_EXCEPTIONS_ROUNDING_SUPPORTED_P
#define TARGET_FLOAT_EXCEPTIONS_ROUNDING_SUPPORTED_P \
@@ -178,6 +178,167 @@ static unsigned HOST_WIDE_INT initial_ix
/* Feature tests against the various architecture variations. */
unsigned char ix86_arch_features[X86_ARCH_LAST];
+struct ix86_target_opts
+{
+ const char *option; /* option string */
+ HOST_WIDE_INT mask; /* isa mask options */
+};
+
+/* This table is ordered so that options like -msse4.2 that imply other
+ ISAs come first. Target string will be displayed in the same order. */
+static struct ix86_target_opts isa2_opts[] =
+{
+ { "-mcx16", OPTION_MASK_ISA_CX16 },
+ { "-mvaes", OPTION_MASK_ISA_VAES },
+ { "-mrdpid", OPTION_MASK_ISA_RDPID },
+ { "-mpconfig", OPTION_MASK_ISA_PCONFIG },
+ { "-mwbnoinvd", OPTION_MASK_ISA_WBNOINVD },
+ { "-mavx512vp2intersect", OPTION_MASK_ISA_AVX512VP2INTERSECT },
+ { "-msgx", OPTION_MASK_ISA_SGX },
+ { "-mavx5124vnniw", OPTION_MASK_ISA_AVX5124VNNIW },
+ { "-mavx5124fmaps", OPTION_MASK_ISA_AVX5124FMAPS },
+ { "-mhle", OPTION_MASK_ISA_HLE },
+ { "-mmovbe", OPTION_MASK_ISA_MOVBE },
+ { "-mclzero", OPTION_MASK_ISA_CLZERO },
+ { "-mmwaitx", OPTION_MASK_ISA_MWAITX },
+ { "-mmovdir64b", OPTION_MASK_ISA_MOVDIR64B },
+ { "-mwaitpkg", OPTION_MASK_ISA_WAITPKG },
+ { "-mcldemote", OPTION_MASK_ISA_CLDEMOTE },
+ { "-mptwrite", OPTION_MASK_ISA_PTWRITE },
+ { "-mavx512bf16", OPTION_MASK_ISA_AVX512BF16 },
+ { "-menqcmd", OPTION_MASK_ISA_ENQCMD }
+};
+static struct ix86_target_opts isa_opts[] =
+{
+ { "-mavx512vpopcntdq", OPTION_MASK_ISA_AVX512VPOPCNTDQ },
+ { "-mavx512bitalg", OPTION_MASK_ISA_AVX512BITALG },
+ { "-mvpclmulqdq", OPTION_MASK_ISA_VPCLMULQDQ },
+ { "-mgfni", OPTION_MASK_ISA_GFNI },
+ { "-mavx512vnni", OPTION_MASK_ISA_AVX512VNNI },
+ { "-mavx512vbmi2", OPTION_MASK_ISA_AVX512VBMI2 },
+ { "-mavx512vbmi", OPTION_MASK_ISA_AVX512VBMI },
+ { "-mavx512ifma", OPTION_MASK_ISA_AVX512IFMA },
+ { "-mavx512vl", OPTION_MASK_ISA_AVX512VL },
+ { "-mavx512bw", OPTION_MASK_ISA_AVX512BW },
+ { "-mavx512dq", OPTION_MASK_ISA_AVX512DQ },
+ { "-mavx512er", OPTION_MASK_ISA_AVX512ER },
+ { "-mavx512pf", OPTION_MASK_ISA_AVX512PF },
+ { "-mavx512cd", OPTION_MASK_ISA_AVX512CD },
+ { "-mavx512f", OPTION_MASK_ISA_AVX512F },
+ { "-mavx2", OPTION_MASK_ISA_AVX2 },
+ { "-mfma", OPTION_MASK_ISA_FMA },
+ { "-mxop", OPTION_MASK_ISA_XOP },
+ { "-mfma4", OPTION_MASK_ISA_FMA4 },
+ { "-mf16c", OPTION_MASK_ISA_F16C },
+ { "-mavx", OPTION_MASK_ISA_AVX },
+/*{ "-msse4" OPTION_MASK_ISA_SSE4 }, */
+ { "-msse4.2", OPTION_MASK_ISA_SSE4_2 },
+ { "-msse4.1", OPTION_MASK_ISA_SSE4_1 },
+ { "-msse4a", OPTION_MASK_ISA_SSE4A },
+ { "-mssse3", OPTION_MASK_ISA_SSSE3 },
+ { "-msse3", OPTION_MASK_ISA_SSE3 },
+ { "-maes", OPTION_MASK_ISA_AES },
+ { "-msha", OPTION_MASK_ISA_SHA },
+ { "-mpclmul", OPTION_MASK_ISA_PCLMUL },
+ { "-msse2", OPTION_MASK_ISA_SSE2 },
+ { "-msse", OPTION_MASK_ISA_SSE },
+ { "-m3dnowa", OPTION_MASK_ISA_3DNOW_A },
+ { "-m3dnow", OPTION_MASK_ISA_3DNOW },
+ { "-mmmx", OPTION_MASK_ISA_MMX },
+ { "-mrtm", OPTION_MASK_ISA_RTM },
+ { "-mprfchw", OPTION_MASK_ISA_PRFCHW },
+ { "-mrdseed", OPTION_MASK_ISA_RDSEED },
+ { "-madx", OPTION_MASK_ISA_ADX },
+ { "-mprefetchwt1", OPTION_MASK_ISA_PREFETCHWT1 },
+ { "-mclflushopt", OPTION_MASK_ISA_CLFLUSHOPT },
+ { "-mxsaves", OPTION_MASK_ISA_XSAVES },
+ { "-mxsavec", OPTION_MASK_ISA_XSAVEC },
+ { "-mxsaveopt", OPTION_MASK_ISA_XSAVEOPT },
+ { "-mxsave", OPTION_MASK_ISA_XSAVE },
+ { "-mabm", OPTION_MASK_ISA_ABM },
+ { "-mbmi", OPTION_MASK_ISA_BMI },
+ { "-mbmi2", OPTION_MASK_ISA_BMI2 },
+ { "-mlzcnt", OPTION_MASK_ISA_LZCNT },
+ { "-mtbm", OPTION_MASK_ISA_TBM },
+ { "-mpopcnt", OPTION_MASK_ISA_POPCNT },
+ { "-msahf", OPTION_MASK_ISA_SAHF },
+ { "-mcrc32", OPTION_MASK_ISA_CRC32 },
+ { "-mfsgsbase", OPTION_MASK_ISA_FSGSBASE },
+ { "-mrdrnd", OPTION_MASK_ISA_RDRND },
+ { "-mpku", OPTION_MASK_ISA_PKU },
+ { "-mlwp", OPTION_MASK_ISA_LWP },
+ { "-mfxsr", OPTION_MASK_ISA_FXSR },
+ { "-mclwb", OPTION_MASK_ISA_CLWB },
+ { "-mshstk", OPTION_MASK_ISA_SHSTK },
+ { "-mmovdiri", OPTION_MASK_ISA_MOVDIRI }
+};
+
+/* Return 1 if TRAIT NAME is present in the OpenMP context's
+ device trait set, return 0 if not present in any OpenMP context in the
+ whole translation unit, or -1 if not present in the current OpenMP context
+ but might be present in another OpenMP context in the same TU. */
+
+int
+ix86_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
+ const char *name)
+{
+ switch (trait)
+ {
+ case omp_device_kind:
+ return strcmp (name, "cpu") == 0;
+ case omp_device_arch:
+ if (strcmp (name, "x86") == 0)
+ return 1;
+ if (TARGET_64BIT)
+ {
+ if (TARGET_X32)
+ return strcmp (name, "x32") == 0;
+ else
+ return strcmp (name, "x86_64") == 0;
+ }
+ if (strcmp (name, "ia32") == 0 || strcmp (name, "i386") == 0)
+ return 1;
+ if (strcmp (name, "i486") == 0)
+ return ix86_arch != PROCESSOR_I386 ? 1 : -1;
+ if (strcmp (name, "i586") == 0)
+ return (ix86_arch != PROCESSOR_I386
+ && ix86_arch != PROCESSOR_I486) ? 1 : -1;
+ if (strcmp (name, "i686") == 0)
+ return (ix86_arch != PROCESSOR_I386
+ && ix86_arch != PROCESSOR_I486
+ && ix86_arch != PROCESSOR_LAKEMONT
+ && ix86_arch != PROCESSOR_PENTIUM) ? 1 : -1;
+ return 0;
+ case omp_device_isa:
+ for (int i = 0; i < 2; i++)
+ {
+ struct ix86_target_opts *opts = i ? isa2_opts : isa_opts;
+ size_t nopts = i ? ARRAY_SIZE (isa2_opts) : ARRAY_SIZE (isa_opts);
+ HOST_WIDE_INT mask = i ? ix86_isa_flags2 : ix86_isa_flags;
+ for (size_t n = 0; n < nopts; n++)
+ {
+ const char *option = opts[n].option + 2;
+ /* -msse4.2 and -msse4.1 options contain dot, which is not valid
+ in identifiers. Use underscore instead, and handle sse4
+ as an alias to sse4_2. */
+ if (opts[n].mask == OPTION_MASK_ISA_SSE4_2)
+ {
+ option = "sse4_2";
+ if (strcmp (name, "sse4") == 0)
+ return (mask & opts[n].mask) != 0 ? 1 : -1;
+ }
+ else if (opts[n].mask == OPTION_MASK_ISA_SSE4_1)
+ option = "sse4_1";
+ if (strcmp (name, option) == 0)
+ return (mask & opts[n].mask) != 0 ? 1 : -1;
+ }
+ }
+ return 0;
+ default:
+ gcc_unreachable ();
+ }
+}
+
/* Return a string that documents the current -m options. The caller is
responsible for freeing the string. */
@@ -187,101 +348,6 @@ ix86_target_string (HOST_WIDE_INT isa, H
const char *arch, const char *tune,
enum fpmath_unit fpmath, bool add_nl_p, bool add_abi_p)
{
- struct ix86_target_opts
- {
- const char *option; /* option string */
- HOST_WIDE_INT mask; /* isa mask options */
- };
-
- /* This table is ordered so that options like -msse4.2 that imply other
- ISAs come first. Target string will be displayed in the same order. */
- static struct ix86_target_opts isa2_opts[] =
- {
- { "-mcx16", OPTION_MASK_ISA_CX16 },
- { "-mvaes", OPTION_MASK_ISA_VAES },
- { "-mrdpid", OPTION_MASK_ISA_RDPID },
- { "-mpconfig", OPTION_MASK_ISA_PCONFIG },
- { "-mwbnoinvd", OPTION_MASK_ISA_WBNOINVD },
- { "-mavx512vp2intersect", OPTION_MASK_ISA_AVX512VP2INTERSECT },
- { "-msgx", OPTION_MASK_ISA_SGX },
- { "-mavx5124vnniw", OPTION_MASK_ISA_AVX5124VNNIW },
- { "-mavx5124fmaps", OPTION_MASK_ISA_AVX5124FMAPS },
- { "-mhle", OPTION_MASK_ISA_HLE },
- { "-mmovbe", OPTION_MASK_ISA_MOVBE },
- { "-mclzero", OPTION_MASK_ISA_CLZERO },
- { "-mmwaitx", OPTION_MASK_ISA_MWAITX },
- { "-mmovdir64b", OPTION_MASK_ISA_MOVDIR64B },
- { "-mwaitpkg", OPTION_MASK_ISA_WAITPKG },
- { "-mcldemote", OPTION_MASK_ISA_CLDEMOTE },
- { "-mptwrite", OPTION_MASK_ISA_PTWRITE },
- { "-mavx512bf16", OPTION_MASK_ISA_AVX512BF16 },
- { "-menqcmd", OPTION_MASK_ISA_ENQCMD }
- };
- static struct ix86_target_opts isa_opts[] =
- {
- { "-mavx512vpopcntdq", OPTION_MASK_ISA_AVX512VPOPCNTDQ },
- { "-mavx512bitalg", OPTION_MASK_ISA_AVX512BITALG },
- { "-mvpclmulqdq", OPTION_MASK_ISA_VPCLMULQDQ },
- { "-mgfni", OPTION_MASK_ISA_GFNI },
- { "-mavx512vnni", OPTION_MASK_ISA_AVX512VNNI },
- { "-mavx512vbmi2", OPTION_MASK_ISA_AVX512VBMI2 },
- { "-mavx512vbmi", OPTION_MASK_ISA_AVX512VBMI },
- { "-mavx512ifma", OPTION_MASK_ISA_AVX512IFMA },
- { "-mavx512vl", OPTION_MASK_ISA_AVX512VL },
- { "-mavx512bw", OPTION_MASK_ISA_AVX512BW },
- { "-mavx512dq", OPTION_MASK_ISA_AVX512DQ },
- { "-mavx512er", OPTION_MASK_ISA_AVX512ER },
- { "-mavx512pf", OPTION_MASK_ISA_AVX512PF },
- { "-mavx512cd", OPTION_MASK_ISA_AVX512CD },
- { "-mavx512f", OPTION_MASK_ISA_AVX512F },
- { "-mavx2", OPTION_MASK_ISA_AVX2 },
- { "-mfma", OPTION_MASK_ISA_FMA },
- { "-mxop", OPTION_MASK_ISA_XOP },
- { "-mfma4", OPTION_MASK_ISA_FMA4 },
- { "-mf16c", OPTION_MASK_ISA_F16C },
- { "-mavx", OPTION_MASK_ISA_AVX },
-/* { "-msse4" OPTION_MASK_ISA_SSE4 }, */
- { "-msse4.2", OPTION_MASK_ISA_SSE4_2 },
- { "-msse4.1", OPTION_MASK_ISA_SSE4_1 },
- { "-msse4a", OPTION_MASK_ISA_SSE4A },
- { "-mssse3", OPTION_MASK_ISA_SSSE3 },
- { "-msse3", OPTION_MASK_ISA_SSE3 },
- { "-maes", OPTION_MASK_ISA_AES },
- { "-msha", OPTION_MASK_ISA_SHA },
- { "-mpclmul", OPTION_MASK_ISA_PCLMUL },
- { "-msse2", OPTION_MASK_ISA_SSE2 },
- { "-msse", OPTION_MASK_ISA_SSE },
- { "-m3dnowa", OPTION_MASK_ISA_3DNOW_A },
- { "-m3dnow", OPTION_MASK_ISA_3DNOW },
- { "-mmmx", OPTION_MASK_ISA_MMX },
- { "-mrtm", OPTION_MASK_ISA_RTM },
- { "-mprfchw", OPTION_MASK_ISA_PRFCHW },
- { "-mrdseed", OPTION_MASK_ISA_RDSEED },
- { "-madx", OPTION_MASK_ISA_ADX },
- { "-mprefetchwt1", OPTION_MASK_ISA_PREFETCHWT1 },
- { "-mclflushopt", OPTION_MASK_ISA_CLFLUSHOPT },
- { "-mxsaves", OPTION_MASK_ISA_XSAVES },
- { "-mxsavec", OPTION_MASK_ISA_XSAVEC },
- { "-mxsaveopt", OPTION_MASK_ISA_XSAVEOPT },
- { "-mxsave", OPTION_MASK_ISA_XSAVE },
- { "-mabm", OPTION_MASK_ISA_ABM },
- { "-mbmi", OPTION_MASK_ISA_BMI },
- { "-mbmi2", OPTION_MASK_ISA_BMI2 },
- { "-mlzcnt", OPTION_MASK_ISA_LZCNT },
- { "-mtbm", OPTION_MASK_ISA_TBM },
- { "-mpopcnt", OPTION_MASK_ISA_POPCNT },
- { "-msahf", OPTION_MASK_ISA_SAHF },
- { "-mcrc32", OPTION_MASK_ISA_CRC32 },
- { "-mfsgsbase", OPTION_MASK_ISA_FSGSBASE },
- { "-mrdrnd", OPTION_MASK_ISA_RDRND },
- { "-mpku", OPTION_MASK_ISA_PKU },
- { "-mlwp", OPTION_MASK_ISA_LWP },
- { "-mfxsr", OPTION_MASK_ISA_FXSR },
- { "-mclwb", OPTION_MASK_ISA_CLWB },
- { "-mshstk", OPTION_MASK_ISA_SHSTK },
- { "-mmovdiri", OPTION_MASK_ISA_MOVDIRI }
- };
-
/* Flag options. */
static struct ix86_target_opts flag_opts[] =
{
@@ -8,3 +8,10 @@ ALL_HOST_OBJS += mkoffload.o
mkoffload$(exeext): mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) $(LIBDEPS)
$(LINKER) $(ALL_LINKERFLAGS) $(LDFLAGS) -o $@ \
mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) $(LIBS)
+
+omp-device-properties: $(srcdir)/config/i386/i386-options.c
+ echo kind: cpu > omp-device-properties
+ echo arch: x86 x86_64 i386 i486 i586 i686 ia32 >> omp-device-properties
+ echo isa: sse4 `sed -n '/^static struct ix86_target_opts isa2\?_opts\[\] =/,/^};/p' \
+ $(srcdir)/config/i386/i386-options.c | \
+ sed -n 's/",.*$$//;s/\./_/;s/^ { "-m//p'` >> omp-device-properties
@@ -10,3 +10,8 @@ mkoffload$(exeext): mkoffload.o collect-
mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) $(LIBS)
MULTILIB_OPTIONS = mgomp
+
+omp-device-properties: $(srcdir)/config/nvptx/nvptx.c
+ echo kind: gpu > omp-device-properties
+ echo arch: nvptx >> omp-device-properties
+ echo isa: sm_30 sm_35 >> omp-device-properties
@@ -5474,6 +5474,32 @@ nvptx_simt_vf ()
return PTX_WARP_SIZE;
}
+/* Return 1 if TRAIT NAME is present in the OpenMP context's
+ device trait set, return 0 if not present in any OpenMP context in the
+ whole translation unit, or -1 if not present in the current OpenMP context
+ but might be present in another OpenMP context in the same TU. */
+
+int
+nvptx_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
+ const char *name)
+{
+ switch (trait)
+ {
+ case omp_device_kind:
+ return strcmp (name, "gpu") == 0;
+ case omp_device_arch:
+ return strcmp (name, "nvptx") == 0;
+ case omp_device_isa:
+ if (strcmp (name, "sm_30") == 0)
+ return !TARGET_SM35;
+ if (strcmp (name, "sm_35") == 0)
+ return TARGET_SM35;
+ return 0;
+ default:
+ gcc_unreachable ();
+ }
+}
+
static bool
nvptx_welformed_vector_length_p (int l)
{
@@ -6539,6 +6565,9 @@ nvptx_set_current_function (tree fndecl)
#undef TARGET_SIMT_VF
#define TARGET_SIMT_VF nvptx_simt_vf
+#undef TARGET_OMP_DEVICE_KIND_ARCH_ISA
+#define TARGET_OMP_DEVICE_KIND_ARCH_ISA nvptx_omp_device_kind_arch_isa
+
#undef TARGET_GOACC_VALIDATE_DIMS
#define TARGET_GOACC_VALIDATE_DIMS nvptx_goacc_validate_dims
@@ -0,0 +1,63 @@
+/* { dg-do compile { target c } } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-additional-options "-mno-sse3" { target { i?86-*-* x86_64-*-* } } } */
+
+#undef i386
+void f01 (void);
+#pragma omp declare variant (f01) match (device={isa(avx512f,avx512bw)})
+void f02 (void);
+void f03 (void);
+#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)})
+void f04 (void);
+void f05 (void);
+#pragma omp declare variant (f05) match (device={kind(gpu)})
+void f06 (void);
+void f07 (void);
+#pragma omp declare variant (f07) match (device={kind(cpu)})
+void f08 (void);
+void f09 (void);
+#pragma omp declare variant (f09) match (device={isa(sm_35)})
+void f10 (void);
+void f11 (void);
+#pragma omp declare variant (f11) match (device={arch(nvptx)})
+void f12 (void);
+void f13 (void);
+#pragma omp declare variant (f13) match (device={arch(i386),isa(sse4)})
+void f14 (void);
+void f15 (void);
+#pragma omp declare variant (f15) match (device={isa(sse4,ssse3),arch(i386)})
+void f16 (void);
+void f17 (void);
+#pragma omp declare variant (f17) match (device={kind(any,fpga)})
+void f18 (void);
+
+void
+test1 (void)
+{
+ int i;
+ f02 (); /* { dg-final { scan-tree-dump-times "f02 \\\(\\\);" 1 "gimple" } } */
+ f14 (); /* { dg-final { scan-tree-dump-times "f14 \\\(\\\);" 1 "gimple" } } */
+ f18 (); /* { dg-final { scan-tree-dump-times "f18 \\\(\\\);" 1 "gimple" } } */
+}
+
+#if defined(__i386__) || defined(__x86_64__)
+__attribute__((target ("avx512f,avx512bw")))
+#endif
+void
+test2 (void)
+{
+ f04 (); /* { dg-final { scan-tree-dump-times "f03 \\\(\\\);" 1 "gimple" { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+ /* { dg-final { scan-tree-dump-times "f04 \\\(\\\);" 1 "gimple" { target { { ! lp64 } || { ! { i?86-*-* x86_64-*-* } } } } } } */
+ f16 (); /* { dg-final { scan-tree-dump-times "f15 \\\(\\\);" 1 "gimple" { target ia32 } } } */
+ /* { dg-final { scan-tree-dump-times "f16 \\\(\\\);" 1 "gimple" { target { ! ia32 } } } } */
+}
+
+void
+test3 (void)
+{
+ f06 (); /* { dg-final { scan-tree-dump-times "f06 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
+ f08 (); /* { dg-final { scan-tree-dump-times "f07 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
+ f10 (); /* { dg-final { scan-tree-dump-times "f10 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
+ f12 (); /* { dg-final { scan-tree-dump-times "f12 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* } } } } } */
+ /* { dg-final { scan-tree-dump-times "f11 \\\(\\\);" 1 "gimple" { target { nvptx*-*-* } } } } */
+}
@@ -0,0 +1,77 @@
+/* { dg-do compile { target c } } */
+/* { dg-additional-options "-foffload=disable -fdump-tree-gimple" } */
+/* { dg-additional-options "-mavx512bw" { target { i?86-*-* x86_64-*-* } } } */
+
+#undef i386
+void f01 (void);
+#pragma omp declare variant (f01) match (device={isa(avx512f,avx512bw)})
+void f02 (void);
+void f03 (void);
+#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)})
+void f04 (void);
+void f05 (void);
+#pragma omp declare variant (f05) match (device={kind(gpu)})
+void f06 (void);
+void f07 (void);
+#pragma omp declare variant (f07) match (device={kind(cpu)})
+void f08 (void);
+void f09 (void);
+#pragma omp declare variant (f09) match (device={isa(sm_35)})
+void f10 (void);
+void f11 (void);
+#pragma omp declare variant (f11) match (device={arch(nvptx)})
+void f12 (void);
+void f13 (void);
+#pragma omp declare variant (f13) match (device={arch(i386),isa(sse4)})
+void f14 (void);
+void f15 (void);
+#pragma omp declare variant (f15) match (device={isa(sse4,ssse3),arch(i386)})
+void f16 (void);
+void f17 (void);
+#pragma omp declare variant (f17) match (device={kind(any,fpga)})
+void f18 (void);
+
+#pragma omp declare target
+void
+test1 (void)
+{
+ int i;
+ f02 (); /* { dg-final { scan-tree-dump-times "f01 \\\(\\\);" 1 "gimple" { target i?86-*-* x86_64-*-* } } } */
+ /* { dg-final { scan-tree-dump-times "f02 \\\(\\\);" 1 "gimple" { target { ! { i?86-*-* x86_64-*-* } } } } } */
+ f14 (); /* { dg-final { scan-tree-dump-times "f13 \\\(\\\);" 1 "gimple" { target ia32 } } } */
+ /* { dg-final { scan-tree-dump-times "f14 \\\(\\\);" 1 "gimple" { target { ! ia32 } } } } */
+ f18 (); /* { dg-final { scan-tree-dump-times "f18 \\\(\\\);" 1 "gimple" } } */
+}
+#pragma omp end declare target
+
+#if defined(__i386__) || defined(__x86_64__)
+__attribute__((target ("avx512f,avx512bw")))
+#endif
+void
+test2 (void)
+{
+ #pragma omp target
+ f04 (); /* { dg-final { scan-tree-dump-times "f03 \\\(\\\);" 1 "gimple" { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+ /* { dg-final { scan-tree-dump-times "f04 \\\(\\\);" 1 "gimple" { target { { ! lp64 } || { ! { i?86-*-* x86_64-*-* } } } } } } */
+ #pragma omp target
+ f16 (); /* { dg-final { scan-tree-dump-times "f15 \\\(\\\);" 1 "gimple" { target ia32 } } } */
+ /* { dg-final { scan-tree-dump-times "f16 \\\(\\\);" 1 "gimple" { target { ! ia32 } } } } */
+}
+
+void
+test3 (void)
+{
+ f06 (); /* { dg-final { scan-tree-dump-times "f06 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
+ f08 (); /* { dg-final { scan-tree-dump-times "f07 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
+}
+#pragma omp declare target to (test3)
+
+void
+test4 (void)
+{
+ #pragma omp target
+ f10 (); /* { dg-final { scan-tree-dump-times "f10 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
+ #pragma omp target
+ f12 (); /* { dg-final { scan-tree-dump-times "f12 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* } } } } } */
+ /* { dg-final { scan-tree-dump-times "f11 \\\(\\\);" 1 "gimple" { target { nvptx*-*-* } } } } */
+}
@@ -811,6 +811,8 @@ LN
LN_S
AWK
SET_MAKE
+omp_device_property_deps
+omp_device_properties
accel_dir_suffix
real_target_noncanonical
enable_as_accelerator
@@ -7879,12 +7881,20 @@ fi
for tgt in `echo $enable_offload_targets | sed 's/,/ /g'`; do
+ tgt_dir=`echo $tgt | sed -n 's/.*=//p'`
tgt=`echo $tgt | sed 's/=.*//'`
if echo "$tgt" | grep "^hsa" > /dev/null ; then
enable_hsa=1
else
enable_offloading=1
+ if test -n "$tgt_dir"; then
+ omp_device_property="${tgt_dir}/lib/gcc/\$(real_target_noncanonical)/\$(version)/accel/${tgt}/omp-device-properties"
+ else
+ omp_device_property="\$(libsubdir)/accel/${tgt}/omp-device-properties"
+ fi
+ omp_device_properties="${omp_device_properties} ${tgt}=${omp_device_property}"
+ omp_device_property_deps="${omp_device_property_deps} ${omp_device_property}"
fi
if test x"$offload_targets" = x; then
@@ -7894,6 +7904,9 @@ for tgt in `echo $enable_offload_targets
fi
done
+
+
+
cat >>confdefs.h <<_ACEOF
#define OFFLOAD_TARGETS "$offload_targets"
_ACEOF
@@ -18851,7 +18864,7 @@ else
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
-#line 18854 "configure"
+#line 18867 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
@@ -18957,7 +18970,7 @@ else
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
-#line 18960 "configure"
+#line 18973 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
@@ -6103,6 +6103,13 @@ to use it.
Return number of threads in SIMT thread group on the target.
@end deftypefn
+@deftypefn {Target Hook} int TARGET_OMP_DEVICE_KIND_ARCH_ISA (enum omp_device_kind_arch_isa @var{trait}, const char *@var{name})
+Return 1 if @var{trait} @var{name} is present in the OpenMP context's
+device trait set, return 0 if not present in any OpenMP context in the
+whole translation unit, or -1 if not present in the current OpenMP context
+but might be present in another OpenMP context in the same TU.
+@end deftypefn
+
@deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree @var{decl}, int *@var{dims}, int @var{fn_level}, unsigned @var{used})
This hook should check the launch dimensions provided for an OpenACC
compute region, or routine. Defaulted values are represented as -1