[RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Jakub Jelinek]

Hi!

The following patch attempts to implement the OpenMP declare variant (and
later on metadirective) device set arch/isa selectors.

The standard makes it implementation defined what an arch is and what is
isa, but I think because there is no selector like target that arch should
mostly contain identifiers that match the ABI incompatible stuff (target,
perhaps whether it is 32-bit or 64-bit, plus endianity where needed etc.)
and keep isa to be identifiers for the ISAs, or perhaps where there are no
clear ISA names say architecture variants or revisions or similar.

I've only implemented i386 and nvptx so far, will leave the rest to
port maintainers; would be nice to coordinate what is added a little bit
with other implementations like LLVM, if they'd be willing to coordinate.

The target hook returns a tri-state value, 0 for doesn't match and will not
match anywhere in the translation unit, 1 for matches and -1 for doesn't
match in the current context, but could match in some other function.
On targets that don't support target attribute or something similar,
returning just 0 or 1 might be enough, -1 is meant for cases where e.g.
during parsing of the pragma when we do not know in which context it will be
called we can signal "don't know whether it will match or not".

The patch doesn't just add a hook, but also infrastructure through which
the --enable-as-accelerator-for= configured compilers can tell the host
compiler what identifiers they do support for kind, arch and isa, so that
for calls used in contexts that are or might be offloaded, we can defer
decisions until after IPA where we know for sure if it is the offloading
version say for nvptx, or offloading version for host fallback, or something
yet different.
Initially I wanted the target hook to handle name == NULL by printing the
list of kind/arch/isa and some undocumented option that the host compiler
would call the accelerator compiler with, but then realized it wouldn't
really work with canadian crosses, so the current version just uses
sed when there are too many values that it needs to be maintained in one
place.

Tested on x86_64-linux with offloading to
nvptx-none,x86_64-intelmicemul-linux-gnu.
Will bootstrap/regtest also without any offloading configured.

Does this approach look reasonable and is it ok with the backend maintainers
listed in To:?  Martin listed for HSA, I'm afraid right now not really sure
at which point it would be possible to distinguish hsa guarded targeted code
from host targeted one.  CCed some backend maintainers for thoughts on what
would be reasonable values for the target hook on their backends.

2019-10-29  Jakub Jelinek  <jakub@redhat.com>

	* configure.ac: Compute and substitute omp_device_properties and
	omp_device_property_deps.
	* Makefile.in (generated_files): Add omp-device-properties.h.
	(omp-general.o): Depend on omp-device-properties.h.
	(omp_device_properties): New make variable.
	(omp-device-properties.h, s-omp-device-properties-h,
	install-omp-device-properties): New goals.
	(install): Depend on install-omp-device-properties for accelerators.
	* target.def (TARGET_OMP_DEVICE_KIND_ARCH_ISA): New target hook.
	* target.h (enum omp_device_kind_arch_isa): New enum.
	* doc/tm.texi.in: Add placeholder for TARGET_OMP_DEVICE_KIND_ARCH_ISA
	documentation.
	* omp-general.c: Include omp-device-properties.h.
	(omp_max_simt_vf): Expect OFFLOAD_TARGET_NAMES to be separated by
	colon instead of comma.
	(omp_offload_device_kind_arch_isa, omp_maybe_offloaded): New
	functions.
	(omp_context_selector_matches): Implement device set arch/isa
	selectors, improve device set kind selector handling.
	* config/i386/i386-options.h (ix86_omp_device_kind_arch_isa): Declare.
	* config/i386/i386.c (TARGET_SIMD_CLONE_ADJUST,
	TARGET_SIMD_CLONE_USABLE): Formatting fix.
	(TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to
	ix86_omp_device_kind_arch_isa.
	* config/i386/i386-options.c (struct ix86_target_opts): Move type
	definition from ix86_target_string to file scope.
	(isa2_opts, isa_opts): Moved arrays from ix86_target_string function
	to file scope.
	(ix86_omp_device_kind_arch_isa): New function.
	(ix86_target_string): Moved struct ix86_target_opts, isa2_opts and
	isa_opts definitions to file scope.
	* config/i386/t-intelmic (omp-device-properties): New goal.
	* config/nvptx/t-nvptx (omp-device-properties): Likewise.
	* config/nvptx/nvptx.c (nvptx_omp_device_kind_arch_isa): New function.
	(TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to
	nvptx_omp_device_kind_arch_isa.
	* configure: Regenerate.
	* doc/tm.texi: Regenerate.
testsuite/
	* c-c++-common/gomp/declare-variant-9.c: New test.
	* c-c++-common/gomp/declare-variant-10.c: New test.

--- gcc/configure.ac.jj	2019-10-29 12:09:19.703952533 +0100
+++ gcc/configure.ac	2019-10-29 12:41:16.337394934 +0100
@@ -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
--- gcc/Makefile.in.jj	2019-10-29 12:09:22.643907176 +0100
+++ gcc/Makefile.in	2019-10-29 12:49:27.389822150 +0100
@@ -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
 
 #\f
 # 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.
--- gcc/target.def.jj	2019-10-29 12:08:05.469097638 +0100
+++ gcc/target.def	2019-10-29 12:16:05.039700360 +0100
@@ -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_"
--- gcc/target.h.jj	2019-10-29 12:08:07.821061357 +0100
+++ gcc/target.h	2019-10-29 12:16:05.027700545 +0100
@@ -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;
--- gcc/doc/tm.texi.in.jj	2019-10-29 12:08:04.390114280 +0100
+++ gcc/doc/tm.texi.in	2019-10-29 12:16:05.038700375 +0100
@@ -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
--- gcc/omp-general.c.jj	2019-10-29 12:09:22.632907346 +0100
+++ gcc/omp-general.c	2019-10-29 16:12:02.062230891 +0100
@@ -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"))
--- gcc/config/i386/i386-options.h.jj	2019-06-10 19:42:14.404796162 +0200
+++ gcc/config/i386/i386-options.h	2019-10-29 12:16:05.004700900 +0100
@@ -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,
--- gcc/config/i386/i386.c.jj	2019-10-29 12:09:23.415895270 +0100
+++ gcc/config/i386/i386.c	2019-10-29 12:16:04.984701208 +0100
@@ -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 \
--- gcc/config/i386/i386-options.c.jj	2019-10-29 12:09:23.414895285 +0100
+++ gcc/config/i386/i386-options.c	2019-10-29 12:16:04.952701702 +0100
@@ -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[] =
   {
--- gcc/config/i386/t-intelmic.jj	2015-04-21 08:39:10.801458081 +0200
+++ gcc/config/i386/t-intelmic	2019-10-29 12:16:05.004700900 +0100
@@ -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
--- gcc/config/nvptx/t-nvptx.jj	2017-04-20 15:00:59.146774473 +0200
+++ gcc/config/nvptx/t-nvptx	2019-10-29 12:16:05.004700900 +0100
@@ -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
--- gcc/config/nvptx/nvptx.c.jj	2019-10-29 12:09:22.691906436 +0100
+++ gcc/config/nvptx/nvptx.c	2019-10-29 12:16:05.005700884 +0100
@@ -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
 
--- gcc/testsuite/c-c++-common/gomp/declare-variant-9.c.jj	2019-10-29 15:07:00.367048135 +0100
+++ gcc/testsuite/c-c++-common/gomp/declare-variant-9.c	2019-10-29 15:10:13.842082046 +0100
@@ -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*-*-* } } } } */
+}
--- gcc/testsuite/c-c++-common/gomp/declare-variant-10.c.jj	2019-10-29 15:12:15.163222136 +0100
+++ gcc/testsuite/c-c++-common/gomp/declare-variant-10.c	2019-10-29 16:22:43.062414098 +0100
@@ -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*-*-* } } } } */
+}
--- gcc/configure.jj	2019-10-29 12:08:07.824061310 +0100
+++ gcc/configure	2019-10-29 12:41:45.081951651 +0100
@@ -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
--- gcc/doc/tm.texi.jj	2019-10-29 12:08:04.389114296 +0100
+++ gcc/doc/tm.texi	2019-10-29 12:16:05.037700391 +0100
@@ -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

	Jakub

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Segher Boessenkool]

Hi!

On Tue, Oct 29, 2019 at 06:15:31PM +0100, Jakub Jelinek wrote:
> The standard makes it implementation defined what an arch is and what is
> isa, but I think because there is no selector like target that arch should
> mostly contain identifiers that match the ABI incompatible stuff (target,
> perhaps whether it is 32-bit or 64-bit, plus endianity where needed etc.)
> and keep isa to be identifiers for the ISAs, or perhaps where there are no
> clear ISA names say architecture variants or revisions or similar.
> 
> I've only implemented i386 and nvptx so far, will leave the rest to
> port maintainers; would be nice to coordinate what is added a little bit
> with other implementations like LLVM, if they'd be willing to coordinate.

What would this be used for?  Can you give some more context?

There already are a lot of different ways to get information about the
execution environment you're running on; why is this any better?


Segher

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Jakub Jelinek]

On Tue, Oct 29, 2019 at 05:40:11PM -0500, Segher Boessenkool wrote:
> On Tue, Oct 29, 2019 at 06:15:31PM +0100, Jakub Jelinek wrote:
> > The standard makes it implementation defined what an arch is and what is
> > isa, but I think because there is no selector like target that arch should
> > mostly contain identifiers that match the ABI incompatible stuff (target,
> > perhaps whether it is 32-bit or 64-bit, plus endianity where needed etc.)
> > and keep isa to be identifiers for the ISAs, or perhaps where there are no
> > clear ISA names say architecture variants or revisions or similar.
> > 
> > I've only implemented i386 and nvptx so far, will leave the rest to
> > port maintainers; would be nice to coordinate what is added a little bit
> > with other implementations like LLVM, if they'd be willing to coordinate.
> 
> What would this be used for?  Can you give some more context?

https://www.openmp.org/spec-html/5.0/openmpse11.html#x41-480002.3
OpenMP 5.0 has two directives, declare variant which I'm implementing right
now and metadirective, which I'll be working on next stage1.
declare variant is a direct call redirection based on context, where one can
provide a specialized function implementation for a particular OpenMP
construct, compiler implementation, CPU architecture, etc.
The metadirective allows specialization of OpenMP pragmas, use this OpenMP pragma
only in some OpenMP context and not in another one, or only on certain
architecture, isa, whatever.
These actually don't query anything at runtime, it is purely compile time
specialization, based on what either the whole translation unit or a
particular function are compiled for.

> There already are a lot of different ways to get information about the
> execution environment you're running on; why is this any better?

There is no question of better or worse, it is simply a part of a standard
that GCC is trying to implement, like we try to implement all of C++20, we
also try to implement all of OpenMP or OpenACC etc.
The exact identifiers are implementation defined though, and that is why we
need to think of what is reasonable for each target (at least each target
where OpenMP is used often, powerpc is certainly one of those).

	Jakub

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Segher Boessenkool]

On Tue, Oct 29, 2019 at 11:57:40PM +0100, Jakub Jelinek wrote:
> On Tue, Oct 29, 2019 at 05:40:11PM -0500, Segher Boessenkool wrote:
> > On Tue, Oct 29, 2019 at 06:15:31PM +0100, Jakub Jelinek wrote:
> > There already are a lot of different ways to get information about the
> > execution environment you're running on; why is this any better?
> 
> There is no question of better or worse, it is simply a part of a standard
> that GCC is trying to implement, like we try to implement all of C++20, we
> also try to implement all of OpenMP or OpenACC etc.

We now get three levels: "arch", "isa", "extensions".  Power has only at
most half of such levels: arch is Power, isa is Power, you can say things
like AltiVec (aka VMX) and VSX are extensions (that's what the X stands
for anyway :-) ).  In older versions of the architecture we had (optional)
"categories", but that is gone as well.

So yes I wonder how we should map reality to these three levels.  And I
also do wonder how this is better than one level as anything else has :-)

> The exact identifiers are implementation defined though, and that is why we
> need to think of what is reasonable for each target (at least each target
> where OpenMP is used often, powerpc is certainly one of those).

Do you know any OpenMP on Power users we can ask for their opinions?


Segher

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Szabolcs Nagy]

On 29/10/2019 17:15, Jakub Jelinek wrote:
> +void f03 (void);
> +#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)})
> +void f04 (void);

1) it's not clear from the omp spec what is the intended
syntax for isa-name, arch-name and extension-name, but
i expected strings in "".

what if an arch or isa name contains ',' ')' etc?

we were planing to use things like

isa("sve")
arch("armv8.2-a+sve")
extension("scalable")

i think we can drop the ", but it looks a bit weird:
normal pp-token parsing of directives would break
the arch name up into multiple tokens (unless it's
special cased like include <...>, at least inside
_Pragma("omp ...") there is no expectation of normal
pp-token parsing), either way is fine with me, but
it may be worth asking clarification from omp?


2) does f03 need to be declared before the declare variant
pragma appears?

for simd variants it means i need to declare the function
with the right simd types and attributes.

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Jakub Jelinek]

On Wed, Oct 30, 2019 at 02:12:30PM +0000, Szabolcs Nagy wrote:
> On 29/10/2019 17:15, Jakub Jelinek wrote:
> > +void f03 (void);
> > +#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)})
> > +void f04 (void);
> 
> 1) it's not clear from the omp spec what is the intended
> syntax for isa-name, arch-name and extension-name, but
> i expected strings in "".

Yes, it is indeed not clear, subject to ongoing discussions.
My reading of the spec was that all the *-name-list are comma
separated lists of identifiers, some others in the committee
want now (yesterday's discussions) string literals instead
when I and others pointed out that isa(core-avx512) can't be valid,
but strangely only for isa/arch/extension but not e.g. for
kind or vendor which would still take identifier lists etc.
My preference at this point would be probably to allow
in all vendor/kind/arch/isa/extension lists of either
identifiers or string literals, so for names which don't contain
characters invalid in identifiers users could choose what to write,
so both isa(avx512f,avx512bw) and isa("avx512f","avx512bw")
would be valid (and so would be isa(avx512f,"avx512bw")), obviously
for something that is not a valid identifier users wouldn't have a choice,
armv8.2-a+sve is not an identifier.

> what if an arch or isa name contains ',' ')' etc?
> 
> we were planing to use things like
> 
> isa("sve")
> arch("armv8.2-a+sve")
> extension("scalable")

That is subject to yet another ongoing discussion in the committee,
what shall be the meaning of isa and what shall be the meaning of arch,
I think we need something that will hold the GCC target or something similar
and something that holds the instruction sets for that target.
extension, given it is in the implementation trait set, is IMNSHO not meant
to hold device extensions, but rather software extensions or something
similar.
ARM is an OpenMP member, so if you want, you can participate too.
https://github.com/OpenMP/spec/issues/2028
is where I'm trying to track all the declare variant issues that need
clarification (plus in two examples tickets).

> 2) does f03 need to be declared before the declare variant
> pragma appears?

Yes.  For C++ the spec says that the actual function declaration
is determined by what would be called at the point of the pragma with
the given id-expression and arguments with types from the function
declaration on which the pragma is used, so it can involve ADL, needs to
deal with function overloading etc.

> for simd variants it means i need to declare the function
> with the right simd types and attributes.

For simd it is actually not finished yet, what needs to be done is that
given the declare simd clauses used as properties of the simd selector
the FEs use some target hook that will guide it how to transform
the parameters like targetm.simd_clone.compute_vecsize_and_simdlen
does and for C tries to just match the types against it and determine
through that the ABI and perhaps missing simd clauses like
notinbranch/inbranch, simdlen etc., for C++ actually for each possibility
will try to construct a call with such arguments and then compare the types.

I'd like to make non-simd working first though, for some cases it already
works and replacement is done in the gimplifier, but scoring needs to be
added, then some way to keep such info in the cgraph (will need to talk to
Honza) and after IPA perform another attempt to redirect.

	Jakub

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Szabolcs Nagy]

On 30/10/2019 14:48, Jakub Jelinek wrote:
> ARM is an OpenMP member, so if you want, you can participate too.
> https://github.com/OpenMP/spec/issues/2028
> is where I'm trying to track all the declare variant issues that need
> clarification (plus in two examples tickets).

it's unfortunate that neither the mailing list nor
the github repo for the spec are public, i'll try
to get access to them to see the discussions.
thanks for the explanations.

>> for simd variants it means i need to declare the function
>> with the right simd types and attributes.
> 
> For simd it is actually not finished yet, what needs to be done is that
> given the declare simd clauses used as properties of the simd selector
> the FEs use some target hook that will guide it how to transform
> the parameters like targetm.simd_clone.compute_vecsize_and_simdlen
> does and for C tries to just match the types against it and determine
> through that the ABI and perhaps missing simd clauses like
> notinbranch/inbranch, simdlen etc., for C++ actually for each possibility
> will try to construct a call with such arguments and then compare the types.

this omp declare variant mechanism seems fairly complicated.

i need is a way to specify a simd variant unambiguously,
which requires is_inbranch, simd_len and vector_call_abi
setting as far as i can tell (potentially a symbol_name
too if the vector abi mangled name is not good enough).

i think i can extend the simd attribute to do this e.g.

__attribute__((simd("notinbranch", 4, "sse2")))
float expf(float);

would declare the _ZGVbN4v_expf simd variant of expf.
(multiple attributes can be used to declare multiple
variants.)

or

__attribute__((simd("notinbranch", 4, "sse2", "my_vexpf")))
float expf(float);

or maybe

typedef float vfloat __attribute__((vector_size(16)));
vfloat my_vexpf(vfloat);

__attribute__((simd("notinbranch", 4, "sse2", my_vexpf)))
float expf(float);

if we allow custom symbol name for the simd variant.

here "sse2" is not specifying a gcc target nor instruction
set, but the vector call convention, e.g. on x86_64 there
could be "sse2", "avx", "avx2" and "avx512", on aarch64
"advsimd" and "sve".

i thought this would match the omp isa, but based on your
description omp isa will be something else.

there is a further complication that vector length agnostic
(scalable) sve calls need a special simd len value: i'd
reserve 0 for it, but it seems internally that means
'simd len is unset' so that has to change.

i will try to prepare an initial patch for such attribute
to make the proposal more concrete, but if you have any
concerns please let me know.

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Richard Sandiford]

Thanks for implementing this.

Jakub Jelinek <jakub@redhat.com> writes:
> On Wed, Oct 30, 2019 at 02:12:30PM +0000, Szabolcs Nagy wrote:
>> On 29/10/2019 17:15, Jakub Jelinek wrote:
>> > +void f03 (void);
>> > +#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)})
>> > +void f04 (void);
>> 
>> 1) it's not clear from the omp spec what is the intended
>> syntax for isa-name, arch-name and extension-name, but
>> i expected strings in "".
>
> Yes, it is indeed not clear, subject to ongoing discussions.
> My reading of the spec was that all the *-name-list are comma
> separated lists of identifiers, some others in the committee
> want now (yesterday's discussions) string literals instead
> when I and others pointed out that isa(core-avx512) can't be valid,
> but strangely only for isa/arch/extension but not e.g. for
> kind or vendor which would still take identifier lists etc.

Might be completely wrong, but wouldn't the identifiers be subject to
macro expansion?  That would make it harder to use the pragmas safely
in system headers.

Richard

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Jakub Jelinek]

On Thu, Oct 31, 2019 at 09:16:00AM +0000, Richard Sandiford wrote:
> > Yes, it is indeed not clear, subject to ongoing discussions.
> > My reading of the spec was that all the *-name-list are comma
> > separated lists of identifiers, some others in the committee
> > want now (yesterday's discussions) string literals instead
> > when I and others pointed out that isa(core-avx512) can't be valid,
> > but strangely only for isa/arch/extension but not e.g. for
> > kind or vendor which would still take identifier lists etc.
> 
> Might be completely wrong, but wouldn't the identifiers be subject to
> macro expansion?  That would make it harder to use the pragmas safely
> in system headers.

Well, all the identifiers in the OpenMP/OpenACC pragmas after the omp
are subject to macro expansion, including clause names etc., so even
#define device foobar
#pragma omp declare variant (foo) match (device={kind(host)})
is a problem, or macro redefinition of declare, variant, match,
kind or host.  That said, for the arch and isa it is perhaps a bigger
problem, as can be seen in the c-c++-common/gomp/declare-variant-9.c
testcase where I had to #undef i386 because in non-strict modes
that is defined...

	Jakub

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Szabolcs Nagy]

On 31/10/2019 09:35, Jakub Jelinek wrote:
> On Thu, Oct 31, 2019 at 09:16:00AM +0000, Richard Sandiford wrote:
>>> Yes, it is indeed not clear, subject to ongoing discussions.
>>> My reading of the spec was that all the *-name-list are comma
>>> separated lists of identifiers, some others in the committee
>>> want now (yesterday's discussions) string literals instead
>>> when I and others pointed out that isa(core-avx512) can't be valid,
>>> but strangely only for isa/arch/extension but not e.g. for
>>> kind or vendor which would still take identifier lists etc.
>>
>> Might be completely wrong, but wouldn't the identifiers be subject to
>> macro expansion?  That would make it harder to use the pragmas safely
>> in system headers.
> 
> Well, all the identifiers in the OpenMP/OpenACC pragmas after the omp
> are subject to macro expansion, including clause names etc., so even
> #define device foobar
> #pragma omp declare variant (foo) match (device={kind(host)})
> is a problem, or macro redefinition of declare, variant, match,
> kind or host.  That said, for the arch and isa it is perhaps a bigger
> problem, as can be seen in the c-c++-common/gomp/declare-variant-9.c
> testcase where I had to #undef i386 because in non-strict modes
> that is defined...

that is surprising, preproc directives are normally not
subject to macro expansion unless stated explicitly so
gcc normally don't macro expand pragmas.. i consider this
a bug in the omp spec and one more reason why _Pragma
should be used instead of #pragma.

Re: [RFC PATCH] targetm.omp.device_kind_arch_isa and OpenMP declare variant kind/arch/isa handling

[Martin Jambor]

Hi,

On Tue, Oct 29 2019, Jakub Jelinek wrote:
> Does this approach look reasonable and is it ok with the backend maintainers
> listed in To:?  Martin listed for HSA, I'm afraid right now not really sure
> at which point it would be possible to distinguish hsa guarded targeted code
> from host targeted one.

Right, since the general approach is to heavily rely on falling back on
the host, I don't think you can do much better than what the patch does.

Thanks,

Martin


> CCed some backend maintainers for thoughts on what
> would be reasonable values for the target hook on their backends.
>
> 2019-10-29  Jakub Jelinek  <jakub@redhat.com>
>
> 	* configure.ac: Compute and substitute omp_device_properties and
> 	omp_device_property_deps.
> 	* Makefile.in (generated_files): Add omp-device-properties.h.
> 	(omp-general.o): Depend on omp-device-properties.h.
> 	(omp_device_properties): New make variable.
> 	(omp-device-properties.h, s-omp-device-properties-h,
> 	install-omp-device-properties): New goals.
> 	(install): Depend on install-omp-device-properties for accelerators.
> 	* target.def (TARGET_OMP_DEVICE_KIND_ARCH_ISA): New target hook.
> 	* target.h (enum omp_device_kind_arch_isa): New enum.
> 	* doc/tm.texi.in: Add placeholder for TARGET_OMP_DEVICE_KIND_ARCH_ISA
> 	documentation.
> 	* omp-general.c: Include omp-device-properties.h.
> 	(omp_max_simt_vf): Expect OFFLOAD_TARGET_NAMES to be separated by
> 	colon instead of comma.
> 	(omp_offload_device_kind_arch_isa, omp_maybe_offloaded): New
> 	functions.
> 	(omp_context_selector_matches): Implement device set arch/isa
> 	selectors, improve device set kind selector handling.
> 	* config/i386/i386-options.h (ix86_omp_device_kind_arch_isa): Declare.
> 	* config/i386/i386.c (TARGET_SIMD_CLONE_ADJUST,
> 	TARGET_SIMD_CLONE_USABLE): Formatting fix.
> 	(TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to
> 	ix86_omp_device_kind_arch_isa.
> 	* config/i386/i386-options.c (struct ix86_target_opts): Move type
> 	definition from ix86_target_string to file scope.
> 	(isa2_opts, isa_opts): Moved arrays from ix86_target_string function
> 	to file scope.
> 	(ix86_omp_device_kind_arch_isa): New function.
> 	(ix86_target_string): Moved struct ix86_target_opts, isa2_opts and
> 	isa_opts definitions to file scope.
> 	* config/i386/t-intelmic (omp-device-properties): New goal.
> 	* config/nvptx/t-nvptx (omp-device-properties): Likewise.
> 	* config/nvptx/nvptx.c (nvptx_omp_device_kind_arch_isa): New function.
> 	(TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to
> 	nvptx_omp_device_kind_arch_isa.
> 	* configure: Regenerate.
> 	* doc/tm.texi: Regenerate.
> testsuite/
> 	* c-c++-common/gomp/declare-variant-9.c: New test.
> 	* c-c++-common/gomp/declare-variant-10.c: New test.
>