diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc
index 854e987dc79..1fd90a1cfb9 100644
--- a/gcc/c-family/c-attribs.cc
+++ b/gcc/c-family/c-attribs.cc
@@ -521,6 +521,8 @@ const struct attribute_spec c_common_gnu_attributes[] =
 			      handle_omp_declare_variant_attribute, NULL },
   { "omp declare variant variant", 0, -1, true,  false, false, false,
 			      handle_omp_declare_variant_attribute, NULL },
+  { "omp declare variant overload", 0, -1, true,  false, false, false,
+			      handle_omp_declare_variant_attribute, NULL },
   { "simd",		      0, 1, true,  false, false, false,
 			      handle_simd_attribute, NULL },
   { "omp declare target",     0, -1, true, false, false, false,
@@ -4020,8 +4022,8 @@ handle_omp_declare_simd_attribute (tree *, tree, tree, int, bool *)
   return NULL_TREE;
 }
 
-/* Handle an "omp declare variant {base,variant}" attribute; arguments as in
-   struct attribute_spec.handler.  */
+/* Handle an "omp declare variant {base,variant,overload}" attribute;
+   arguments as in struct attribute_spec.handler.  */
 
 static tree
 handle_omp_declare_variant_attribute (tree *, tree, tree, int, bool *)
diff --git a/gcc/coverage.cc b/gcc/coverage.cc
index ad55f0f1909..0dc156d9c05 100644
--- a/gcc/coverage.cc
+++ b/gcc/coverage.cc
@@ -646,6 +646,7 @@ coverage_begin_function (unsigned lineno_checksum, unsigned cfg_checksum)
 		     (DECL_ASSEMBLER_NAME (current_function_decl)));
   gcov_write_unsigned (DECL_ARTIFICIAL (current_function_decl)
 		       && !DECL_FUNCTION_VERSIONED (current_function_decl)
+		       && !DECL_FUNCTION_OMP_VARIANT (current_function_decl)
 		       && !DECL_LAMBDA_FUNCTION_P (current_function_decl));
   gcov_write_filename (remap_profile_filename (startloc.file));
   gcov_write_unsigned (startloc.line);
diff --git a/gcc/cp/call.cc b/gcc/cp/call.cc
index c7efc5b077a..44848342b25 100644
--- a/gcc/cp/call.cc
+++ b/gcc/cp/call.cc
@@ -13093,6 +13093,29 @@ joust (struct z_candidate *cand1, struct z_candidate *cand2, bool warn,
 	}
     }
 
+  if (TREE_CODE (cand1->fn) == FUNCTION_DECL
+      && DECL_FUNCTION_OMP_VARIANT (cand1->fn)
+      && TREE_CODE (cand2->fn) == FUNCTION_DECL
+      && DECL_FUNCTION_OMP_VARIANT (cand2->fn))
+    {
+      tree f1 = TREE_TYPE (cand1->fn);
+      tree f2 = TREE_TYPE (cand2->fn);
+      tree p1 = TYPE_ARG_TYPES (f1);
+      tree p2 = TYPE_ARG_TYPES (f2);
+      if (compparms (p1, p2)
+	  && same_type_p (TREE_TYPE (f1), TREE_TYPE (f2)))
+	{
+	  tree a1 = lookup_attribute ("omp declare variant overload",
+				      DECL_ATTRIBUTES (cand1->fn));
+	  tree a2 = lookup_attribute ("omp declare variant overload",
+				      DECL_ATTRIBUTES (cand2->fn));
+	  if (!a1 && a2)
+	    return 1;
+	  else if (a1 && !a2)
+	    return -1;
+	}
+    }
+
   /* If the two function declarations represent the same function (this can
      happen with declarations in multiple scopes and arg-dependent lookup),
      arbitrarily choose one.  But first make sure the default args we're
diff --git a/gcc/cp/class.cc b/gcc/cp/class.cc
index 6fdb56abfb9..d9751a28014 100644
--- a/gcc/cp/class.cc
+++ b/gcc/cp/class.cc
@@ -1180,6 +1180,18 @@ add_method (tree type, tree method, bool via_using)
 	  && maybe_version_functions (method, fn, true))
 	continue;
 
+      if (flag_openmp
+	  && (lookup_attribute ("omp declare variant overload",
+				DECL_ATTRIBUTES (method))
+	      || lookup_attribute ("omp declare variant overload",
+				   DECL_ATTRIBUTES (fn))))
+	{
+	  bool record = (!DECL_FUNCTION_OMP_VARIANT (method)
+			 || !DECL_FUNCTION_OMP_VARIANT (fn));
+	  maybe_omp_variant_functions (method, fn, record);
+	  continue;
+	}
+
       if (DECL_INHERITED_CTOR (method))
 	{
 	  if (!DECL_INHERITED_CTOR (fn))
@@ -8759,6 +8771,9 @@ resolve_address_of_overloaded_function (tree target_type,
 	mark_versions_used (fn);
     }
 
+  // TBD.
+  gcc_assert (!DECL_FUNCTION_OMP_VARIANT (fn));
+
   /* If we're doing overload resolution purely for the purpose of
      determining conversion sequences, we should not consider the
      function used.  If this conversion sequence is selected, the
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index cb89d372b23..c32ac60c2a9 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -1845,6 +1845,11 @@ struct GTY(()) cp_omp_begin_assumes_data {
   bool attr_syntax;
 };
 
+struct GTY(()) cp_omp_declare_variant_attr {
+  bool attr_syntax;
+  tree selector;
+};
+
 /* Global state.  */
 
 struct GTY(()) saved_scope {
@@ -1894,6 +1899,7 @@ struct GTY(()) saved_scope {
   hash_map<tree, tree> *GTY((skip)) x_local_specializations;
   vec<cp_omp_declare_target_attr, va_gc> *omp_declare_target_attribute;
   vec<cp_omp_begin_assumes_data, va_gc> *omp_begin_assumes;
+  vec<cp_omp_declare_variant_attr, va_gc> *omp_declare_variant_attribute;
 
   struct saved_scope *prev;
 };
@@ -6937,6 +6943,7 @@ extern bool member_like_constrained_friend_p	(tree);
 extern bool fns_correspond			(tree, tree);
 extern int decls_match				(tree, tree, bool = true);
 extern bool maybe_version_functions		(tree, tree, bool);
+extern void maybe_omp_variant_functions		(tree, tree, bool);
 extern bool merge_default_template_args		(tree, tree, bool);
 extern tree duplicate_decls			(tree, tree,
 						 bool hiding = false,
diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc
index 1844c923b7f..4181f42d648 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -1180,6 +1180,21 @@ decls_match (tree newdecl, tree olddecl, bool record_versions /* = true */)
 				      || !DECL_FUNCTION_VERSIONED (olddecl)));
 	  return 0;
 	}
+      if (flag_openmp
+	  && types_match
+	  && (lookup_attribute ("omp declare variant overload",
+				DECL_ATTRIBUTES (newdecl))
+	      || lookup_attribute ("omp declare variant overload",
+				   DECL_ATTRIBUTES (olddecl))))
+	{
+	  if (record_versions)
+	    {
+	      bool record = (!DECL_FUNCTION_OMP_VARIANT (newdecl)
+			     || !DECL_FUNCTION_OMP_VARIANT (olddecl));
+	      maybe_omp_variant_functions (newdecl, olddecl, record);
+	    }
+	  return 0;
+	}
     }
   else if (TREE_CODE (newdecl) == TEMPLATE_DECL)
     {
@@ -1243,6 +1258,17 @@ maybe_mark_function_versioned (tree decl)
     }
 }
 
+static void
+maybe_mark_omp_variant_function (tree decl)
+{
+  if (!DECL_FUNCTION_OMP_VARIANT (decl))
+    {
+      DECL_FUNCTION_OMP_VARIANT (decl) = 1;
+      if (DECL_ASSEMBLER_NAME_SET_P (decl))
+	mangle_decl (decl);
+    }
+}
+
 /* NEWDECL and OLDDECL have identical signatures.  If they are
    different versions adjust them and return true.
    If RECORD is set to true, record function versions.  */
@@ -1277,6 +1303,116 @@ maybe_version_functions (tree newdecl, tree olddecl, bool record)
   return true;
 }
 
+void
+maybe_omp_variant_functions (tree newdecl, tree olddecl, bool record)
+{
+  maybe_mark_omp_variant_function (olddecl);
+  if (DECL_LOCAL_DECL_P (olddecl))
+    {
+      /* Is this meaningful for these functions?  */
+      olddecl = DECL_LOCAL_DECL_ALIAS (olddecl);
+      maybe_mark_omp_variant_function (olddecl);
+    }
+
+  maybe_mark_omp_variant_function (newdecl);
+  if (DECL_LOCAL_DECL_P (newdecl))
+    {
+      /* again?  */
+      if (!DECL_LOCAL_DECL_ALIAS (newdecl))
+	{
+	  if (!DECL_LOCAL_DECL_ALIAS (newdecl))
+	    push_local_extern_decl_alias (newdecl);
+	  newdecl = DECL_LOCAL_DECL_ALIAS (newdecl);
+	  maybe_mark_omp_variant_function (newdecl);
+	}
+    }
+
+  if (record)
+    cgraph_node::record_function_versions (olddecl, newdecl);
+
+  if (record)
+    {
+      cgraph_node *oldn = cgraph_node::get (olddecl);
+      cgraph_node *newn = cgraph_node::get (newdecl);
+      cgraph_function_version_info *oldv = oldn->function_version ();
+      cgraph_function_version_info *first_ver, *ver;
+
+      for (ver = oldv; ver; ver = ver->prev)
+	first_ver = ver;
+
+      tree have_base = NULL_TREE;
+      hash_set<tree> linked_variants;
+
+      const char *base_attr_name = "omp declare variant base";
+      size_t base_attr_len = strlen (base_attr_name);
+
+      for (ver = first_ver; ver; ver = ver->next)
+	{
+	  tree fn = ver->this_node->decl;
+	  tree variant = lookup_attribute ("omp declare variant overload",
+					   DECL_ATTRIBUTES (fn));
+	  if (!variant)
+	    {
+	      /* There should only be one base function (non-variant) in the
+		 list.  FIXME: Might not be true for versioned functions.  */
+	      gcc_assert (!have_base);
+	      have_base = fn;
+	      tree attrs = DECL_ATTRIBUTES (fn);
+	      while (attrs)
+		{
+		  tree attr = get_attribute_name (attrs);
+		  size_t ident_len = IDENTIFIER_LENGTH (attr);
+		  if (cmp_attribs (base_attr_name, base_attr_len,
+				   IDENTIFIER_POINTER (attr), ident_len))
+		    {
+		      /* This is a TREE_LIST node used as a tuple of:
+			 (variant decl, context selector, location).  */
+		      tree info = TREE_VALUE (attr);
+		      tree linked_variant = TREE_PURPOSE (info);
+		      /* Linked variant should appear only once.  */
+		      gcc_assert (!linked_variants.contains (linked_variant));
+		      linked_variants.add (linked_variant);
+		    }
+		  attrs = TREE_CHAIN (attrs);
+		}
+	    }
+	}
+
+      /* We haven't seen the base function yet; do nothing.  */
+      if (!have_base)
+	return;
+
+      for (ver = first_ver; ver; ver = ver->next)
+	{
+	  tree fn = ver->this_node->decl;
+	  if (fn == have_base)
+	    continue;
+	  if (linked_variants.contains (fn))
+	    continue;
+	  /* The base function doesn't have a link to this variant yet.  Add
+	     one now.  */
+	  tree attrib = lookup_attribute ("omp declare variant overload",
+					  DECL_ATTRIBUTES (fn));
+	  gcc_assert (attrib);
+	  tree ctx_selector = TREE_VALUE (attrib);
+	  /* This location stuff is nonsense.  FIXME.  */
+	  location_t fn_loc = DECL_SOURCE_LOCATION (fn);
+	  tree wrapped_loc = maybe_wrap_with_location (integer_zero_node,
+						       fn_loc);
+	  cp_id_kind idk = CP_ID_KIND_NONE;
+	  wrapped_loc
+	    = tree_cons (wrapped_loc, build_int_cst (integer_type_node, idk),
+			 build_tree_list (wrapped_loc, integer_zero_node));
+	  tree info = tree_cons (fn, ctx_selector, wrapped_loc);
+	  DECL_ATTRIBUTES (have_base)
+	    = tree_cons (get_identifier (base_attr_name), info,
+			 DECL_ATTRIBUTES (have_base));
+	  if (processing_template_decl)
+	    ATTR_IS_DEPENDENT (DECL_ATTRIBUTES (have_base)) = 1;
+	}
+    }
+}
+
 /* If NEWDECL is `static' and an `extern' was seen previously,
    warn about it.  OLDDECL is the previous declaration.
 
@@ -1547,10 +1683,14 @@ duplicate_function_template_decls (tree newdecl, tree olddecl)
   tree oldres = DECL_TEMPLATE_RESULT (olddecl);
   /* Function template declarations can be differentiated by parameter
      and return type.  */
-  if (compparms (TYPE_ARG_TYPES (TREE_TYPE (oldres)),
-		 TYPE_ARG_TYPES (TREE_TYPE (newres)))
-       && same_type_p (TREE_TYPE (TREE_TYPE (newdecl)),
-		       TREE_TYPE (TREE_TYPE (olddecl))))
+  if ((!lookup_attribute ("omp declare variant overload",
+			  DECL_ATTRIBUTES (newres))
+       && !lookup_attribute ("omp declare variant overload",
+			     DECL_ATTRIBUTES (oldres)))
+      && compparms (TYPE_ARG_TYPES (TREE_TYPE (oldres)),
+		    TYPE_ARG_TYPES (TREE_TYPE (newres)))
+      && same_type_p (TREE_TYPE (TREE_TYPE (newdecl)),
+		      TREE_TYPE (TREE_TYPE (olddecl))))
     {
       /* ... and also by their template-heads and requires-clauses.  */
       if (template_heads_equivalent_p (newdecl, olddecl)
@@ -1970,6 +2110,8 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden)
 	     are not ambiguous.  */
 	  else if ((!DECL_FUNCTION_VERSIONED (newdecl)
 		    && !DECL_FUNCTION_VERSIONED (olddecl))
+		   && (!DECL_FUNCTION_OMP_VARIANT (newdecl)
+		       && !DECL_FUNCTION_OMP_VARIANT (olddecl))
 		   /* Let constrained hidden friends coexist for now, we'll
 		      check satisfaction later.  */
 		   && !member_like_constrained_friend_p (newdecl)
@@ -3143,6 +3285,15 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden)
       cgraph_node::delete_function_version_by_decl (newdecl);
     }
 
+  if (TREE_CODE (newdecl) == FUNCTION_DECL
+      && DECL_FUNCTION_OMP_VARIANT (olddecl))
+    {
+      /* Set the flag for newdecl so that it gets copied to olddecl.  */
+      DECL_FUNCTION_OMP_VARIANT (newdecl) = 1;
+      /* FIXME: Do we need this here?  */
+      //cgraph_node::delete_function_version_by_decl (newdecl);
+    }
+
   if (TREE_CODE (newdecl) == FUNCTION_DECL)
     {
       int function_size;
@@ -8168,6 +8319,20 @@ omp_declare_variant_finalize_one (tree decl, tree attr)
   if (idk == CP_ID_KIND_QUALIFIED)
     variant = finish_call_expr (variant, &args, /*disallow_virtual=*/true,
 				koenig_p, tf_warning_or_error);
+  else if (idk == CP_ID_KIND_NONE
+	   && DECL_NONSTATIC_MEMBER_FUNCTION_P (variant)
+	   && CLASS_TYPE_P (DECL_CONTEXT (decl)))
+    {
+      tree saved_ccp = current_class_ptr;
+      tree saved_ccr = current_class_ref;
+      current_class_ptr = NULL_TREE;
+      current_class_ref = NULL_TREE;
+      inject_this_parameter (DECL_CONTEXT (decl), TYPE_UNQUALIFIED);
+      variant = finish_call_expr (variant, &args, /*disallow_virtual=*/false,
+				  koenig_p, tf_warning_or_error);
+      current_class_ptr = saved_ccp;
+      current_class_ref = saved_ccr;
+    }
   else
     variant = finish_call_expr (variant, &args, /*disallow_virtual=*/false,
 				koenig_p, tf_warning_or_error);
@@ -8194,6 +8359,9 @@ omp_declare_variant_finalize_one (tree decl, tree attr)
 	  error_at (varid_loc, "variant %qD is a built-in", variant);
 	  return true;
 	}
+      else if (lookup_attribute ("omp declare variant overload",
+				 DECL_ATTRIBUTES (variant)))
+	TREE_PURPOSE (TREE_VALUE (attr)) = variant;
       else
 	{
 	  tree construct
diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc
index bee84879023..df0ef66d14a 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -1774,6 +1774,17 @@ cplus_decl_attributes (tree *decl, tree attributes, int flags)
 	}
     }
 
+  if (vec_safe_length (scope_chain->omp_declare_variant_attribute)
+      && TREE_CODE (*decl) == FUNCTION_DECL)
+    {
+      int length = scope_chain->omp_declare_variant_attribute->length ();
+      cp_omp_declare_variant_attr a
+	= (*scope_chain->omp_declare_variant_attribute)[length - 1];
+      tree ctx = copy_list (a.selector);
+      attributes = tree_cons (get_identifier ("omp declare variant overload"),
+			      ctx, attributes);
+    }
+
   tree late_attrs = NULL_TREE;
   if (processing_template_decl)
     {
diff --git a/gcc/cp/lex.cc b/gcc/cp/lex.cc
index 64bcfb18196..672be7e3616 100644
--- a/gcc/cp/lex.cc
+++ b/gcc/cp/lex.cc
@@ -33,6 +33,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "gcc-rich-location.h"
 #include "cp-name-hint.h"
 #include "langhooks.h"
+#include "omp-general.h"
 
 static int interface_strcmp (const char *);
 static void init_cp_pragma (void);
@@ -328,6 +329,8 @@ cxx_init (void)
   init_cp_semantics ();
   init_operators ();
   init_method ();
+  if (flag_openmp)
+    omp_init_mangle ();
 
   current_function_decl = NULL;
 
diff --git a/gcc/cp/mangle.cc b/gcc/cp/mangle.cc
index 0684f0e6038..073d79c2644 100644
--- a/gcc/cp/mangle.cc
+++ b/gcc/cp/mangle.cc
@@ -55,6 +55,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "stor-layout.h"
 #include "flags.h"
 #include "attribs.h"
+#include "omp-general.h"
 
 /* Debugging support.  */
 
@@ -828,6 +829,21 @@ write_mangled_name (const tree decl, bool top_level)
       else
 	gcc_unreachable ();
     }
+
+  if (flag_openmp
+      && TREE_CODE (decl) == FUNCTION_DECL
+      && DECL_FUNCTION_OMP_VARIANT (decl))
+    {
+      tree ctx;
+      if ((ctx = lookup_attribute ("omp declare variant overload",
+				   DECL_ATTRIBUTES (decl))))
+	{
+	  write_string (JOIN_STR "ompvariant");
+	  const char *append
+	    = omp_mangle_context_selector (JOIN_STR, TREE_VALUE (ctx));
+	  write_string (append);
+	}
+    }
 }
 
 /* Returns true if the return type of DECL is part of its signature, and
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 65aeb83758b..95f7381bab2 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -48440,6 +48440,62 @@ cp_parser_omp_begin (cp_parser *parser, cp_token *pragma_tok)
 	    = { in_omp_attribute_pragma, device_type, indirect };
 	  vec_safe_push (scope_chain->omp_declare_target_attribute, a);
 	}
+      else if (strcmp (p, "variant") == 0)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  const char *clause = "";
+	  matching_parens parens;
+	  location_t match_loc = cp_lexer_peek_token (parser->lexer)->location;
+	  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+	    {
+	      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+	      clause = IDENTIFIER_POINTER (id);
+	    }
+	  if (strcmp (clause, "match") != 0)
+	    {
+	      cp_parser_error (parser, "expected %<match%>");
+	      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+	      return;
+	    }
+
+	  cp_lexer_consume_token (parser->lexer);
+
+	  if (!parens.require_open (parser))
+	    {
+	      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+	      return;
+	    }
+
+	  tree ctx = cp_parser_omp_context_selector_specification (parser,
+								   true);
+	  if (ctx == error_mark_node)
+	    {
+	      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+	      return;
+	    }
+
+	  if (vec_safe_length (scope_chain->omp_declare_variant_attribute) > 0)
+	    {
+	      int length
+		= scope_chain->omp_declare_variant_attribute->length ();
+	      cp_omp_declare_variant_attr a
+		= (*scope_chain->omp_declare_variant_attribute)[length - 1];
+	      tree outer_ctx = a.selector;
+	      ctx = omp_merge_context_selectors (outer_ctx, ctx);
+	    }
+
+	  ctx = omp_check_context_selector (match_loc, ctx);
+
+	  if (ctx != error_mark_node)
+	    {
+	      cp_omp_declare_variant_attr a
+		= { parser->lexer->in_omp_attribute_pragma, ctx };
+	      vec_safe_push (scope_chain->omp_declare_variant_attribute, a);
+	    }
+
+	  parens.require_close (parser);
+	  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+	}
       else
 	{
 	  cp_parser_error (parser, "expected %<target%>");
@@ -48486,41 +48542,70 @@ cp_parser_omp_end (cp_parser *parser, cp_token *pragma_tok)
 	  p = IDENTIFIER_POINTER (id);
 	}
       if (strcmp (p, "target") == 0)
-	cp_lexer_consume_token (parser->lexer);
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  cp_parser_require_pragma_eol (parser, pragma_tok);
+	  if (!vec_safe_length (scope_chain->omp_declare_target_attribute))
+	    error_at (pragma_tok->location,
+		      "%<#pragma omp end declare target%> without "
+		      "corresponding %<#pragma omp declare target%> or "
+		      "%<#pragma omp begin declare target%>");
+	  else
+	    {
+	      cp_omp_declare_target_attr
+		a = scope_chain->omp_declare_target_attribute->pop ();
+	      if (a.attr_syntax != in_omp_attribute_pragma)
+		{
+		  if (a.attr_syntax)
+		    error_at (pragma_tok->location,
+			      "%qs in attribute syntax terminated "
+			      "with %qs in pragma syntax",
+			      a.device_type >= 0 ? "begin declare target"
+						 : "declare target",
+			      "end declare target");
+		  else
+		    error_at (pragma_tok->location,
+			      "%qs in pragma syntax terminated "
+			      "with %qs in attribute syntax",
+			      a.device_type >= 0 ? "begin declare target"
+						 : "declare target",
+			      "end declare target");
+		}
+	    }
+	}
+      else if (strcmp (p, "variant") == 0)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  cp_parser_require_pragma_eol (parser, pragma_tok);
+	  if (!vec_safe_length (scope_chain->omp_declare_variant_attribute))
+	    error_at (pragma_tok->location,
+		      "%<#pragma omp end declare variant%> without "
+		      "corresponding %<#pragma omp begin declare variant%>");
+	  else
+	    {
+	      cp_omp_declare_variant_attr
+		a = scope_chain->omp_declare_variant_attribute->pop ();
+	      if (a.attr_syntax != in_omp_attribute_pragma)
+		{
+		  if (a.attr_syntax)
+		    error_at (pragma_tok->location,
+			      "%<begin declare variant%> in attribute syntax "
+			      "terminated with %<end declare variant%> in "
+			      "pragma syntax");
+		  else
+		    error_at (pragma_tok->location,
+			      "%<begin declare variant%> in pragma syntax "
+			      "terminated with %<end declare variant%> in "
+			      "attribute syntax");
+		}
+	    }
+	}
       else
 	{
 	  cp_parser_error (parser, "expected %<target%>");
 	  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
 	  return;
 	}
-      cp_parser_require_pragma_eol (parser, pragma_tok);
-      if (!vec_safe_length (scope_chain->omp_declare_target_attribute))
-	error_at (pragma_tok->location,
-		  "%<#pragma omp end declare target%> without corresponding "
-		  "%<#pragma omp declare target%> or "
-		  "%<#pragma omp begin declare target%>");
-      else
-	{
-	  cp_omp_declare_target_attr
-	    a = scope_chain->omp_declare_target_attribute->pop ();
-	  if (a.attr_syntax != in_omp_attribute_pragma)
-	    {
-	      if (a.attr_syntax)
-		error_at (pragma_tok->location,
-			  "%qs in attribute syntax terminated "
-			  "with %qs in pragma syntax",
-			  a.device_type >= 0 ? "begin declare target"
-					     : "declare target",
-			  "end declare target");
-	      else
-		error_at (pragma_tok->location,
-			  "%qs in pragma syntax terminated "
-			  "with %qs in attribute syntax",
-			  a.device_type >= 0 ? "begin declare target"
-					     : "declare target",
-			  "end declare target");
-	    }
-	}
     }
   else if (strcmp (p, "assumes") == 0)
     {
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 42d020b105d..eb3672dd774 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -1669,6 +1669,18 @@ spec_hasher::equal (spec_entry *e1, spec_entry *e2)
       tree c2 = e2->spec ? get_constraints (e2->spec) : NULL_TREE;
       equal = equivalent_constraints (c1, c2);
     }
+  if (equal && flag_openmp)
+    {
+      tree r1 = DECL_TEMPLATE_RESULT (e1->tmpl);
+      tree r2 = DECL_TEMPLATE_RESULT (e2->tmpl);
+      if (TREE_CODE (r1) == FUNCTION_DECL
+	  && TREE_CODE (r2) == FUNCTION_DECL
+	  && (lookup_attribute ("omp declare variant overload",
+			       DECL_ATTRIBUTES (r1))
+	      || lookup_attribute ("omp declare variant overload",
+				   DECL_ATTRIBUTES (r2))))
+	equal = false;
+    }
   --processing_template_decl;
   --comparing_dependent_aliases;
   --comparing_specializations;
@@ -11775,6 +11787,10 @@ tsubst_contract_attributes (tree decl, tree args, tsubst_flags_t complain, tree
   DECL_ATTRIBUTES (decl) = list;
 }
 
+static tree
+tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
+		      tree lambda_fntype, bool use_spec_table = true);
+
 /* Instantiate a single dependent attribute T (a TREE_LIST), and return either
    T or a new TREE_LIST, possibly a chain in the case of a pack expansion.  */
 
@@ -11812,9 +11828,16 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
 	   && is_attribute_p ("omp declare variant base",
 			      get_attribute_name (t)))
     {
-      ++cp_unevaluated_operand;
-      tree varid = tsubst_expr (TREE_PURPOSE (val), args, complain, in_decl);
-      --cp_unevaluated_operand;
+      tree varid;
+      if (TREE_CODE (TREE_PURPOSE (val)) == FUNCTION_DECL)
+	varid = tsubst_function_decl (TREE_PURPOSE (val), args, complain,
+				      NULL_TREE);
+      else
+	{
+	  ++cp_unevaluated_operand;
+	  varid = tsubst_expr (TREE_PURPOSE (val), args, complain, in_decl);
+	  --cp_unevaluated_operand;
+	}
       tree chain = TREE_CHAIN (val);
       location_t match_loc = cp_expr_loc_or_input_loc (TREE_PURPOSE (chain));
       tree ctx = copy_list (TREE_VALUE (val));
@@ -11908,6 +11931,10 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
 	}
       val = tree_cons (varid, ctx, chain);
     }
+  else if (flag_openmp
+	   && is_attribute_p ("omp declare variant overload",
+			      get_attribute_name (t)))
+    ;
   /* If the first attribute argument is an identifier, don't
      pass it through tsubst.  Attributes like mode, format,
      cleanup and several target specific attributes expect it
@@ -14329,7 +14356,7 @@ maybe_rebuild_function_decl_type (tree decl)
 
 static tree
 tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
-		      tree lambda_fntype, bool use_spec_table = true)
+		      tree lambda_fntype, bool use_spec_table /*= true*/)
 {
   tree gen_tmpl = NULL_TREE, argvec = NULL_TREE;
   hashval_t hash = 0;
@@ -14630,7 +14657,11 @@ tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
   if (flag_openmp)
     if (tree attr = lookup_attribute ("omp declare variant base",
 				      DECL_ATTRIBUTES (r)))
-      omp_declare_variant_finalize (r, attr);
+      {
+	DECL_ATTRIBUTES (r) = tsubst_attributes (DECL_ATTRIBUTES (r), args,
+						 complain, r);
+	omp_declare_variant_finalize (r, attr);
+      }
 
   return r;
 }
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 6634acfda3f..e70feb85ab5 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -3467,6 +3467,13 @@ finish_translation_unit (void)
 	       "#pragma omp end declare target");
       vec_safe_truncate (scope_chain->omp_declare_target_attribute, 0);
     }
+  if (vec_safe_length (scope_chain->omp_declare_variant_attribute))
+    {
+      if (!errorcount)
+	error ("%<omp begin declare variant%> without corresponding "
+	       "%<omp end declare variant%>");
+      vec_safe_truncate (scope_chain->omp_declare_variant_attribute, 0);
+    }
   if (vec_safe_length (scope_chain->omp_begin_assumes))
     {
       if (!errorcount)
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 8281ec67e00..45c68fa72bc 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -10099,6 +10099,8 @@ expand_omp_target (struct omp_region *region)
 	    = DECL_FUNCTION_SPECIFIC_TARGET (current_function_decl);
 	  DECL_FUNCTION_VERSIONED (child_fn2)
 	    = DECL_FUNCTION_VERSIONED (current_function_decl);
+	  DECL_FUNCTION_OMP_VARIANT (child_fn2)
+	    = DECL_FUNCTION_OMP_VARIANT (current_function_decl);
 
 	  fn2_node = cgraph_node::get_create (child_fn2);
 	  fn2_node->offloadable = 1;
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index 361630bcaed..e908adde0b4 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -45,6 +45,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "data-streamer.h"
 #include "streamer-hooks.h"
 #include "opts.h"
+#include "tree-hash-traits.h"
 
 enum omp_requires omp_requires_mask;
 
@@ -1125,6 +1126,8 @@ const char *omp_tss_map[] =
 /* Arrays of property candidates must be null-terminated.  */
 static const char *const kind_properties[] =
   { "host", "nohost", "cpu", "gpu", "fpga", "any", NULL };
+static const char *const kind_abbrevs[] =
+  { "h", "n", "c", "g", "f", "" };
 static const char *const vendor_properties[] =
   { "amd", "arm", "bsc", "cray", "fujitsu", "gnu", "hpe", "ibm", "intel",
     "llvm", "nvidia", "pgi", "ti", "unknown", NULL };
@@ -1132,111 +1135,112 @@ static const char *const extension_properties[] =
   { NULL };
 static const char *const atomic_default_mem_order_properties[] =
   { "seq_cst", "relaxed", "acq_rel", "acquire", "release", NULL };
+static const char *const atomic_default_mem_order_abbrevs[] =
+  { "sc", "rx", "ar", "ac", "re" };
 
 struct omp_ts_info omp_ts_map[] =
   {
-   { "kind",
+   { "kind", "k",
      (1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
      OMP_TRAIT_PROPERTY_NAME_LIST, false,
-     kind_properties
+     kind_properties, kind_abbrevs
    },
-   { "isa",
+   { "isa", "i",
      (1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
      OMP_TRAIT_PROPERTY_NAME_LIST, false,
-     NULL
+     NULL, NULL
    },
-   { "arch",
+   { "arch", "a",
      (1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
      OMP_TRAIT_PROPERTY_NAME_LIST, false,
-     NULL
+     NULL, NULL
    },
-   { "device_num",
+   { "device_num", NULL,
      (1 << OMP_TRAIT_SET_TARGET_DEVICE),
      OMP_TRAIT_PROPERTY_EXPR, false,
-     NULL
+     NULL, NULL
    },
-   { "vendor",
+   { "vendor", "v",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_NAME_LIST, true,
-     vendor_properties,
+     vendor_properties, NULL
    },
-   { "extension",
+   { "extension", "e",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_NAME_LIST, true,
-     extension_properties,
+     extension_properties, NULL
    },
-   { "atomic_default_mem_order",
+   { "atomic_default_mem_order", "a",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_ID, true,
      atomic_default_mem_order_properties,
+     atomic_default_mem_order_abbrevs
    },
-   { "requires",
+   { "requires", "r",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_CLAUSE_LIST, true,
-     NULL
+     NULL, NULL
    },
-   { "unified_address",
+   { "unified_address", "una",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_NONE, true,
-     NULL
+     NULL, NULL
    },
-   { "unified_shared_memory",
+   { "unified_shared_memory", "usm",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_NONE, true,
-     NULL
+     NULL, NULL
    },
-   { "dynamic_allocators",
+   { "dynamic_allocators", "dna",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_NONE, true,
-     NULL
+     NULL, NULL
    },
-   { "reverse_offload",
+   { "reverse_offload", "rvo",
      (1 << OMP_TRAIT_SET_IMPLEMENTATION),
      OMP_TRAIT_PROPERTY_NONE, true,
-     NULL
+     NULL, NULL
    },
-   { "condition",
+   { "condition", NULL,
      (1 << OMP_TRAIT_SET_USER),
      OMP_TRAIT_PROPERTY_EXPR, true,
-     NULL
+     NULL, NULL
    },
-   { "target",
+   { "target", "ta",
      (1 << OMP_TRAIT_SET_CONSTRUCT),
      OMP_TRAIT_PROPERTY_NONE, false,
-     NULL
+     NULL, NULL
    },
-   { "teams",
+   { "teams", "te",
      (1 << OMP_TRAIT_SET_CONSTRUCT),
      OMP_TRAIT_PROPERTY_NONE, false,
-     NULL
+     NULL, NULL
    },
-   { "parallel",
+   { "parallel", "pa",
      (1 << OMP_TRAIT_SET_CONSTRUCT),
      OMP_TRAIT_PROPERTY_NONE, false,
-     NULL
+     NULL, NULL
    },
-   { "for",
+   { "for", "fo",
      (1 << OMP_TRAIT_SET_CONSTRUCT),
      OMP_TRAIT_PROPERTY_NONE, false,
-     NULL
+     NULL, NULL
    },
-   { "simd",
+   { "simd", "si",
      (1 << OMP_TRAIT_SET_CONSTRUCT),
      OMP_TRAIT_PROPERTY_CLAUSE_LIST,  false,
-     NULL
+     NULL, NULL
    },
-   { NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL }  /* OMP_TRAIT_LAST */
+   /* OMP_TRAIT_LAST */
+   { NULL, NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL, NULL }
   };
 
+/* Helper function for omp_context_name_list_prop.  Return string from either
+   an identifier or a string constant.  */
 
-/* Return a name from PROP, a property in selectors accepting
-   name lists.  */
-
-const char *
-omp_context_name_list_prop (tree prop)
+static const char *
+omp_context_name_list_prop_1 (tree val)
 {
-  gcc_assert (OMP_TP_NAME (prop) == OMP_TP_NAMELIST_NODE);
-  tree val = OMP_TP_VALUE (prop);
   switch (TREE_CODE (val))
     {
     case IDENTIFIER_NODE:
@@ -1254,6 +1258,16 @@ omp_context_name_list_prop (tree prop)
     }
 }
 
+/* Return a name from PROP, a property in selectors accepting
+   name lists.  */
+
+const char *
+omp_context_name_list_prop (tree prop)
+{
+  gcc_assert (OMP_TP_NAME (prop) == OMP_TP_NAMELIST_NODE);
+  return omp_context_name_list_prop_1 (OMP_TP_VALUE (prop));
+}
+
 /* Diagnose errors in an OpenMP context selector, return CTX if
    it is correct or error_mark_node otherwise.  */
 
@@ -1797,6 +1811,611 @@ omp_context_selector_matches (tree ctx)
   return ret;
 }
 
+static void
+omp_copy_trait_set (tree from_ts, tree *to_ts)
+{
+  for (tree ts = from_ts; ts; ts = TREE_CHAIN (ts))
+    *to_ts = make_trait_selector (OMP_TS_CODE (ts), OMP_TS_SCORE (ts),
+				  OMP_TS_PROPERTIES (ts), *to_ts);
+}
+
+/* Trait properties are stored in a tree list, and for namelist properties,
+   may use either an IDENTIFIER_POINTER or STRING_CST for their value.  We
+   want these to be considered equal if they represent the same string (e.g.
+   quoted or not in the source program).  Arrange a hash table to allow us to
+   do this.  */
+
+typedef std::pair<tree, tree> omp_trait_prop;
+
+struct omp_trait_prop_hash : typed_noop_remove <omp_trait_prop>
+{
+  typedef omp_trait_prop value_type;
+  typedef omp_trait_prop compare_type;
+
+  static inline hashval_t hash (omp_trait_prop);
+  static const bool empty_zero_p = true;
+  static inline bool is_empty (omp_trait_prop);
+  static inline bool is_deleted (omp_trait_prop);
+  static inline bool equal (const omp_trait_prop &, const omp_trait_prop &);
+  static inline void mark_empty (omp_trait_prop &);
+};
+
+inline hashval_t
+omp_trait_prop_hash::hash (omp_trait_prop t)
+{
+  if (t.first == OMP_TP_NAMELIST_NODE)
+    return htab_hash_string (omp_context_name_list_prop_1 (t.second));
+  if (t.first && t.second)
+    return iterative_hash_expr (t.second, iterative_hash_expr (t.first, 0));
+  else if (t.first)
+    return iterative_hash_expr (t.first, 0);
+  else if (t.second)
+    return iterative_hash_expr (t.second, 0);
+  else
+    gcc_unreachable ();
+}
+
+inline bool
+omp_trait_prop_hash::is_empty (omp_trait_prop t)
+{
+  return !t.first && !t.second;
+}
+
+inline bool
+omp_trait_prop_hash::is_deleted (omp_trait_prop)
+{
+  return false;
+}
+
+inline bool
+omp_trait_prop_hash::equal (const omp_trait_prop &a, const omp_trait_prop &b)
+{
+  if (a.first == OMP_TP_NAMELIST_NODE
+      && b.first == OMP_TP_NAMELIST_NODE)
+    {
+      const char *a_name = omp_context_name_list_prop_1 (a.second);
+      const char *b_name = omp_context_name_list_prop_1 (b.second);
+      return strcmp (a_name, b_name) == 0;
+    }
+
+  if (a.first && a.second && b.first && b.second)
+    return operand_equal_p (a.first, b.first)
+	   && operand_equal_p (a.second, b.second);
+  else if (a.first && b.first)
+    return !a.second && !b.second && a.first == b.first;
+  else if (a.second && b.second)
+    return !a.first && !b.first && operand_equal_p (a.second, b.second);
+  else
+    return false;
+}
+
+inline void
+omp_trait_prop_hash::mark_empty (omp_trait_prop &t)
+{
+  t.first = NULL_TREE;
+  t.second = NULL_TREE;
+}
+
+typedef hash_set<omp_trait_prop_hash> omp_trait_prop_set;
+
+static void
+omp_gather_trait_sets (omp_trait_prop_set sets[], tree scores[], tree inlist)
+{
+  for (tree ts = inlist; ts; ts = TREE_CHAIN (ts))
+    {
+      enum omp_ts_code ts_code = OMP_TS_CODE (ts);
+
+      if (scores && omp_ts_map[ts_code].allow_score)
+	scores[ts_code] = OMP_TS_SCORE (ts);
+
+      for (tree tp = OMP_TS_PROPERTIES (ts); tp; tp = TREE_CHAIN (tp))
+	{
+	  tree name = OMP_TP_NAME (tp);
+	  tree value = OMP_TP_VALUE (tp);
+	  sets[ts_code].add (std::make_pair (name, value));
+	}
+    }
+}
+
+/* We might have things like:
+
+     outer: arch(avx2,sse4.1,3dnow)
+     inner: arch(sse4.1)
+
+     outer: kind(any)  kind(gpu,cpu)  kind(gpu,cpu)
+     inner: kind(gpu)  kind(gpu)      kind(any)
+
+   We want to add the intersection of the inner and outer sets -- most of the
+   time that just means adding the inner set, but there are some special cases
+   that need handling (e.g. "any" on inner, but some more restrictive kind on
+   outer).  */
+
+static void
+omp_combine_trait_properties (omp_trait_prop_set &outer_props,
+			      omp_trait_prop_set &inner_props,
+			      omp_tss_code tss_code, omp_ts_code ts_code,
+			      tree *to_list, tree score)
+{
+  tree anyt = get_identifier ("any");
+  tree combined_traits = NULL_TREE;
+
+  std::pair<tree, tree> any = std::make_pair (OMP_TP_NAMELIST_NODE, anyt);
+
+  if (outer_props.is_empty ())
+    for (std::pair<tree, tree> e : inner_props)
+      combined_traits = make_trait_property (e.first, e.second,
+					     combined_traits);
+  else if (inner_props.is_empty ()
+	   || (tss_code == OMP_TRAIT_SET_DEVICE
+	       && ts_code == OMP_TRAIT_DEVICE_KIND
+	       && inner_props.contains (any)))
+    for (std::pair<tree, tree> e : outer_props)
+      combined_traits = make_trait_property (e.first, e.second,
+					     combined_traits);
+  else
+    for (std::pair<tree, tree> e : inner_props)
+      if (tss_code == OMP_TRAIT_SET_IMPLEMENTATION
+	  && ts_code == OMP_TRAIT_IMPLEMENTATION_ADMO
+	  && !outer_props.contains (e))
+	error ("cannot combine nested %<atomic_default_mem_order%> properties");
+      else if (outer_props.contains (e)
+	       || (tss_code == OMP_TRAIT_SET_DEVICE
+		   && ts_code == OMP_TRAIT_DEVICE_KIND
+		   && outer_props.contains (any)))
+	combined_traits = make_trait_property (e.first, e.second,
+					       combined_traits);
+
+  if (combined_traits)
+    *to_list = make_trait_selector (ts_code, score,
+				    nreverse (combined_traits), *to_list);
+}
+
+static void
+omp_combine_trait_sets (tree outer_ts, tree inner_ts, omp_tss_code tss_code,
+			tree *to_list)
+{
+  unsigned HOST_WIDE_INT used_traits = 0;
+
+  for (tree t = outer_ts; t; t = TREE_CHAIN (t))
+    used_traits |= 1 << OMP_TS_CODE (t);
+  for (tree t = inner_ts; t; t = TREE_CHAIN (t))
+    used_traits |= 1 << OMP_TS_CODE (t);
+
+  omp_trait_prop_set outer_sets[OMP_TRAIT_LAST];
+  omp_trait_prop_set inner_sets[OMP_TRAIT_LAST];
+
+  tree outer_scores[OMP_TRAIT_LAST] = { };
+  tree inner_scores[OMP_TRAIT_LAST] = { };
+
+  omp_gather_trait_sets (outer_sets, outer_scores, outer_ts);
+  omp_gather_trait_sets (inner_sets, inner_scores, inner_ts);
+
+  for (unsigned i = 0; i < OMP_TRAIT_LAST; i++)
+    if (used_traits & (1 << i))
+      {
+	omp_ts_code ts_code = static_cast<omp_ts_code>(i);
+	/* We can technically have a score on both the inner and outer context
+	   selectors (e.g. for the "implementation" trait set selector).  Just
+	   ignore the outer one.  */
+	tree use_score = inner_scores[i] ? inner_scores[i] : outer_scores[i];
+
+	switch (omp_ts_map[ts_code].tp_type)
+	  {
+	  case OMP_TRAIT_PROPERTY_ID:
+	  case OMP_TRAIT_PROPERTY_NAME_LIST:
+	    omp_combine_trait_properties (outer_sets[i], inner_sets[i],
+					  tss_code, ts_code, to_list,
+					  use_score);
+	    break;
+
+	  case OMP_TRAIT_PROPERTY_EXPR:
+	    /* The user trait set selector just contains the 'condition' trait.
+	       No attempt is made to combine conditions in outer/inner context
+	       selectors here -- instead, the outer condition is just
+	       dropped.  */
+	    if (ts_code == OMP_TRAIT_USER_CONDITION && inner_ts && outer_ts)
+	      warning (OPT_Wopenmp, "ignoring condition on enclosing context "
+		       "selector");
+
+	    if (inner_ts)
+	      omp_copy_trait_set (inner_ts, to_list);
+	    else if (outer_ts)
+	      omp_copy_trait_set (outer_ts, to_list);
+	    break;
+
+	  case OMP_TRAIT_PROPERTY_NONE:
+	    switch (ts_code)
+	      {
+	      case OMP_TRAIT_IMPLEMENTATION_UNIFIED_ADDRESS:
+	      case OMP_TRAIT_IMPLEMENTATION_UNIFIED_SHARED_MEMORY:
+	      case OMP_TRAIT_IMPLEMENTATION_DYNAMIC_ALLOCATORS:
+	      case OMP_TRAIT_IMPLEMENTATION_REVERSE_OFFLOAD:
+		/* For these ones, we want the union of the features listed on
+		   the outer and inner context selectors.  We already checked
+		   the trait has been mentioned in one or the other selector
+		   via 'used_traits', so add it here.  */
+		*to_list = make_trait_selector (ts_code, NULL_TREE,
+						NULL_TREE, *to_list);
+		break;
+
+	      default:
+		sorry ("don't know how to merge nested context selectors");
+	      }
+	    break;
+
+	  default:
+	    gcc_unreachable ();
+	  }
+      }
+}
+
+extern tree omp_get_context_selector_list (tree ctx, enum omp_tss_code set);
+
+tree
+omp_merge_context_selectors (tree outer_ctx, tree inner_ctx)
+{
+  tree merged_ctx = NULL_TREE;
+
+  for (unsigned i = OMP_TRAIT_SET_CONSTRUCT; i != OMP_TRAIT_SET_LAST; i++)
+    {
+      omp_tss_code tss_code = static_cast<omp_tss_code>(i);
+      tree outer_ts = omp_get_context_selector_list (outer_ctx, tss_code);
+      tree inner_ts = omp_get_context_selector_list (inner_ctx, tss_code);
+      tree merged_ts = NULL_TREE;
+
+      switch (tss_code)
+	{
+	case OMP_TRAIT_SET_CONSTRUCT:
+	  if (outer_ts)
+	    omp_copy_trait_set (outer_ts, &merged_ts);
+	  if (inner_ts)
+	    omp_copy_trait_set (inner_ts, &merged_ts);
+	  break;
+
+	case OMP_TRAIT_SET_DEVICE:
+	case OMP_TRAIT_SET_TARGET_DEVICE:
+	case OMP_TRAIT_SET_IMPLEMENTATION:
+	case OMP_TRAIT_SET_USER:
+	  if (outer_ts && inner_ts)
+	    omp_combine_trait_sets (outer_ts, inner_ts, tss_code, &merged_ts);
+	  else if (outer_ts)
+	    omp_copy_trait_set (outer_ts, &merged_ts);
+	  else if (inner_ts)
+	    omp_copy_trait_set (inner_ts, &merged_ts);
+	  break;
+
+	default:
+	  gcc_unreachable ();
+	}
+
+      if (merged_ts)
+	merged_ctx = make_trait_set_selector (tss_code, nreverse (merged_ts),
+					      merged_ctx);
+    }
+
+  return nreverse (merged_ctx);
+}
+
+/* OpenMP "declare variant" mangling obstack bits.  These are essentially
+   copied from cp/mangle.cc, but we'll need them for C also, so they need to go
+   here.  */
+
+static struct obstack *omp_mangle_obstack;
+static struct obstack omp_name_obstack;
+static void *omp_name_base;
+
+/* Append a single character to the end of the mangled
+   representation.  */
+#define write_char(CHAR)						\
+  obstack_1grow (omp_mangle_obstack, (CHAR))
+
+/* Append a NUL-terminated string to the end of the mangled
+   representation.  */
+#define write_string(STRING)						\
+  obstack_grow (omp_mangle_obstack, (STRING), strlen (STRING))
+
+static void
+omp_start_mangling (void)
+{
+  obstack_free (&omp_name_obstack, omp_name_base);
+  omp_mangle_obstack = &omp_name_obstack;
+  omp_name_base = obstack_alloc (&omp_name_obstack, 0);
+}
+
+static const char *
+omp_finish_mangling (void)
+{
+  write_char ('\0');
+  return (const char *) obstack_finish (omp_mangle_obstack);
+}
+
+void
+omp_init_mangle (void)
+{
+  gcc_obstack_init (&omp_name_obstack);
+  omp_name_base = obstack_alloc (&omp_name_obstack, 0);
+}
+
+static int
+omp_string_compare (const void *a, const void *b)
+{
+  const char *as = *(char * const *) a;
+  const char *bs = *(char * const *) b;
+  return strcmp (as, bs);
+}
+
+static void
+omp_stringify_sorted_property_set (const char *sep, omp_ts_code ts_code,
+				   omp_trait_prop_set &props)
+{
+  vec<const char *> names = vNULL;
+  hash_map<nofree_string_hash, const char *> abbrevs;
+  omp_tp_type tp_type = omp_ts_map[ts_code].tp_type;
+
+  /* This function only really works for name lists.  */
+  gcc_assert (tp_type == OMP_TRAIT_PROPERTY_NAME_LIST
+	      || tp_type == OMP_TRAIT_PROPERTY_ID);
+
+  if (omp_ts_map[ts_code].prop_abbrevs)
+    for (int i = 0; omp_ts_map[ts_code].valid_properties[i]; i++)
+      abbrevs.put (omp_ts_map[ts_code].valid_properties[i],
+		   omp_ts_map[ts_code].prop_abbrevs[i]);
+
+  for (auto prop : props)
+    {
+      if (prop.first == OMP_TP_NAMELIST_NODE)
+	{
+	  const char *name = omp_context_name_list_prop_1 (prop.second);
+	  if (omp_ts_map[ts_code].prop_abbrevs)
+	    {
+	      const char **repl = abbrevs.get (name);
+	      if (repl)
+		names.safe_push (*repl);
+	      else
+		error ("missing abbreviation for %qs", name);
+	    }
+	  else
+	    names.safe_push (name);
+	}
+      else if (tp_type == OMP_TRAIT_PROPERTY_ID
+	       && TREE_CODE (prop.first) == IDENTIFIER_NODE)
+	{
+	  const char *name = IDENTIFIER_POINTER (prop.first);
+	  if (omp_ts_map[ts_code].prop_abbrevs)
+	    {
+	      const char **repl = abbrevs.get (name);
+	      if (repl)
+		names.safe_push (*repl);
+	      else
+		error ("missing abbreviation for %qs", name);
+	    }
+	  else
+	    names.safe_push (name);
+	}
+      else
+	gcc_unreachable ();
+    }
+
+  names.qsort (omp_string_compare);
+
+  const char *use_sep = "";
+
+  for (auto &n : names)
+    {
+      bool needs_escaping = false;
+      for (int i = 0; n[i]; i++)
+	if (!ISALPHA (n[i]) && !ISDIGIT (n[i]) && n[i] != '_')
+	  needs_escaping = true;
+      if (needs_escaping)
+	{
+	  write_string (use_sep);
+	  for (int i = 0; n[i]; i++)
+	    {
+	      /* We may have an arbirary string constants representing e.g. an
+		 ISA, which could contain characters not valid for symbols in
+		 assembly output.  The below isn't really safe for arbitrary
+		 character encodings (either in source or assembly output), but
+		 should be a good-enough approximation for our purposes.  */
+	      if (ISALPHA (n[i]) || ISDIGIT (n[i]) || n[i] == '_')
+		write_char (n[i]);
+	      else
+		{
+		  char tmpbuf[4];
+		  sprintf (tmpbuf, "%.2x", (unsigned char) n[i]);
+		  write_string (sep);
+		  write_string (tmpbuf);
+		}
+	    }
+	}
+      else
+	{
+	  write_string (use_sep);
+	  write_string (n);
+	}
+      /* If we're using abbreviations for these trait properties, we don't need
+	 extra separators between them.  */
+      if (!omp_ts_map[ts_code].prop_abbrevs)
+	use_sep = sep;
+    }
+}
+
+/* Key:
+
+   C: start construct
+      ta: target
+      te: teams
+      pa: parallel
+      fo: for/do
+      si: simd
+   D: start device
+      i: isa
+	(ordered) [<string> SEP] list
+      a: arch
+	(ordered) [<string> SEP] list
+      k: kind
+	h: host
+	n: nohost
+	c: cpu
+	g: gpu
+	f: fpga
+	(omitted): any
+   T: start target_device
+      (as above, also)
+      N: (ordered) [<device number> SEP] list
+   I: start implementation
+      v: vendor
+	(ordered) [<string> SEP] list
+      e: extension
+	(nothing here)
+      r: requires
+	(tbd)
+      a: atomic_default_mem_order
+	 ar: acq_rel
+	 ac: acquire
+	 rx: relaxed
+	 re: release
+	 sc: seq_cst
+   U: start user
+     (no further encoding, just 'U' or not)
+   SEP: separate trait selector sets
+*/
+
+const char *
+omp_mangle_context_selector (const char *sep, tree ctx)
+{
+  gcc_assert (ctx);
+
+  omp_start_mangling ();
+
+  for (int i = OMP_TRAIT_SET_CONSTRUCT; i != OMP_TRAIT_SET_LAST; i++)
+    {
+      omp_tss_code tss_code = static_cast<omp_tss_code>(i);
+      tree ts_list = omp_get_context_selector_list (ctx, tss_code);
+
+      if (!ts_list)
+	continue;
+
+      switch (tss_code)
+	{
+	case OMP_TRAIT_SET_CONSTRUCT:
+	  {
+	    write_string (sep);
+	    write_char ('C');
+	    for (tree ts = ts_list; ts; ts = TREE_CHAIN (ts))
+	      write_string (omp_ts_map[OMP_TS_CODE (ts)].abbrev_name);
+	  }
+	  break;
+	case OMP_TRAIT_SET_DEVICE:
+	  {
+	    omp_trait_prop_set trait_sets[OMP_TRAIT_LAST];
+
+	    omp_gather_trait_sets (trait_sets, NULL, ts_list);
+	    write_string (sep);
+	    write_char ('D');
+
+	    const char *use_sep = "";
+
+	    for (unsigned j = 0; j < OMP_TRAIT_LAST; j++)
+	      {
+		omp_ts_code ts_code = static_cast<omp_ts_code>(j);
+
+		if (!(omp_ts_map[ts_code].tss_mask
+		      & (1 << OMP_TRAIT_SET_DEVICE)))
+		  continue;
+
+		if (trait_sets[ts_code].is_empty ())
+		  continue;
+
+		write_string (use_sep);
+		write_string (omp_ts_map[ts_code].abbrev_name);
+
+		omp_stringify_sorted_property_set (sep, ts_code,
+						   trait_sets[ts_code]);
+
+		use_sep = sep;
+	      }
+	  }
+	  break;
+	case OMP_TRAIT_SET_TARGET_DEVICE:
+	  {
+	    write_string (sep);
+	    write_char ('T');
+	    /* This should share the 'device' code above, but for now is
+	       unimplemented.  */
+	    gcc_unreachable ();
+	  }
+	  break;
+	case OMP_TRAIT_SET_IMPLEMENTATION:
+	  {
+	    omp_trait_prop_set trait_sets[OMP_TRAIT_LAST];
+
+	    omp_gather_trait_sets (trait_sets, NULL, ts_list);
+	    write_string (sep);
+	    write_char ('I');
+	    const char *use_sep = "";
+
+	    for (unsigned j = 0; j < OMP_TRAIT_LAST; j++)
+	      {
+		omp_ts_code ts_code = static_cast<omp_ts_code>(j);
+
+		if (!(omp_ts_map[ts_code].tss_mask
+		      & (1 << OMP_TRAIT_SET_IMPLEMENTATION)))
+		  continue;
+
+		switch (omp_ts_map[ts_code].tp_type)
+		  {
+		  case OMP_TRAIT_PROPERTY_NONE:
+		    if (omp_get_context_selector (ctx, tss_code, ts_code))
+		      {
+			write_string (use_sep);
+			write_string (omp_ts_map[ts_code].abbrev_name);
+		      }
+		    break;
+
+		  case OMP_TRAIT_PROPERTY_ID:
+		  case OMP_TRAIT_PROPERTY_NAME_LIST:
+		    if (trait_sets[i].is_empty ())
+		      continue;
+
+		    write_string (use_sep);
+		    write_string (omp_ts_map[ts_code].abbrev_name);
+
+		    omp_stringify_sorted_property_set (sep, ts_code,
+						       trait_sets[ts_code]);
+		    break;
+
+		  case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
+		    if (omp_get_context_selector (ctx, tss_code, ts_code))
+		      {
+			write_string (use_sep);
+			write_string ("?????");
+		      }
+		    break;
+
+		  default:
+		    gcc_unreachable ();
+		  }
+
+		use_sep = sep;
+	      }
+	    continue;
+	  }
+	  break;
+	case OMP_TRAIT_SET_USER:
+	  {
+	    write_string (sep);
+	    write_char ('U');
+	    continue;
+	  }
+	  break;
+	default:
+	  gcc_unreachable ();
+	}
+    }
+
+  return omp_finish_mangling ();
+}
+
 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
    in omp_context_selector_set_compare.  */
 
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 88a0b28f483..d9f83f40116 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -170,6 +170,9 @@ extern tree omp_check_context_selector (location_t loc, tree ctx);
 extern void omp_mark_declare_variant (location_t loc, tree variant,
 				      tree construct);
 extern int omp_context_selector_matches (tree);
+extern tree omp_merge_context_selectors (tree, tree);
+extern void omp_init_mangle (void);
+extern const char * omp_mangle_context_selector (const char *, tree);
 extern int omp_context_selector_set_compare (enum omp_tss_code, tree, tree);
 extern tree omp_get_context_selector (tree, enum omp_tss_code,
 				      enum omp_ts_code);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index dd802ca37a6..9993a2e0ee1 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -2089,6 +2089,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
     = DECL_FUNCTION_SPECIFIC_TARGET (current_function_decl);
   DECL_FUNCTION_VERSIONED (decl)
     = DECL_FUNCTION_VERSIONED (current_function_decl);
+  DECL_FUNCTION_OMP_VARIANT (decl)
+    = DECL_FUNCTION_OMP_VARIANT (current_function_decl);
 
   if (omp_maybe_offloaded_ctx (ctx))
     {
diff --git a/gcc/omp-selectors.h b/gcc/omp-selectors.h
index 825a082c939..65011bb610c 100644
--- a/gcc/omp-selectors.h
+++ b/gcc/omp-selectors.h
@@ -80,10 +80,12 @@ extern const char *omp_tss_map [];
    null-terminated array of strings.  */
 struct omp_ts_info {
   const char *name;
+  const char *abbrev_name;
   unsigned int tss_mask;
   enum omp_tp_type tp_type;
   bool allow_score;
   const char * const *valid_properties;
+  const char * const *prop_abbrevs;
 };
 extern struct omp_ts_info omp_ts_map[];
 
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-1.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-1.C
new file mode 100644
index 00000000000..2af1001d6cb
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-1.C
@@ -0,0 +1,82 @@
+#if 1
+#ifdef _OPENMP
+namespace bar {
+#pragma omp begin declare variant match(construct={target})
+#pragma omp begin declare variant match(construct={parallel})
+namespace qux {
+int foo (char *buf, int len)
+{
+  return 5;
+}
+} // namespace qux
+#pragma omp end declare variant
+#pragma omp end declare variant
+} // namespace bar
+#endif
+
+namespace bar {
+namespace qux {
+int foo (char *buf, int len)
+{
+  return 3;
+}
+} // namespace qux
+} // namespace bar
+
+#pragma omp begin declare variant match(implementation={vendor(gnu)})
+#pragma omp begin declare variant match(implementation={vendor(gnu,nvidia)})
+int quux (int c, int d)
+{
+  return c+d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={arch("mips","riscv")})
+#pragma omp begin declare variant match(device={arch("mips"),isa("mips4")})
+int quuux (int c, int d)
+{
+  return c+d;
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(device={arch("riscv"),isa("rv3.5")})
+int quuux (int c, int d)
+{
+  return c+d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int quux (int c, int d)
+{
+  return c+d;
+}
+#endif
+
+#if 1
+#pragma omp begin declare variant match(device={kind(any)})
+#pragma omp begin declare variant match(device={kind(gpu,nohost,fpga)})
+int warble (int c, int d)
+{
+  return c*d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+#endif
+
+#if 1
+#pragma omp begin declare variant match(device={kind(gpu,nohost,fpga)})
+#pragma omp begin declare variant match(device={kind(any)})
+int warble2 (int c, int d)
+{
+  return c*d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+#endif
+
+
+#pragma omp begin declare variant match(device={arch("mips",riscv)})
+#pragma omp begin declare variant match(device={arch(mips,"riscv")})
+#pragma omp end declare variant
+#pragma omp end declare variant
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-11.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-11.C
new file mode 100644
index 00000000000..837eaab1d29
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-11.C
@@ -0,0 +1,20 @@
+#pragma omp begin declare variant match(implementation={dynamic_allocators,unified_shared_memory})
+/* This is a "sorry" for now.  */
+#pragma omp begin declare variant match(implementation={requires(reverse_offload)})
+int foo (int c)
+{
+  return c+5;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int foo (int c)
+{
+  return c+3;
+}
+
+int main(void)
+{
+  foo (6);
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-12.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-12.C
new file mode 100644
index 00000000000..26f0213195b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-12.C
@@ -0,0 +1,32 @@
+#include <stdio.h>
+
+int r_foo (int c, int d)
+{
+  return (c+d) * 3;
+}
+
+#pragma omp declare variant (r_foo) match(implementation={requires(unified_shared_memory)})
+int foo (int c, int d)
+{
+  return c+d;
+}
+
+int main()
+{
+  printf ("host: %d\n", foo (3, 4));
+
+#pragma omp parallel
+  {
+#pragma omp single
+    {
+      printf ("parallel: %d\n", foo (3, 4));
+    }
+  }
+
+#pragma omp target
+  {
+    printf ("target: %d\n", foo (3, 4));
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-2.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-2.C
new file mode 100644
index 00000000000..98d7dff9963
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-2.C
@@ -0,0 +1,26 @@
+#pragma omp begin declare variant match(implementation={vendor(score(5):nvidia)})
+#pragma omp begin declare variant match(implementation={vendor(nvidia)})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(nvidia)})
+#pragma omp begin declare variant match(implementation={vendor(score(7):nvidia)})
+int bar (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(score(9):nvidia)})
+#pragma omp begin declare variant match(implementation={vendor(score(13):"nvidia")})
+int qux (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+  const bool enable_p = true;
+  const bool enable2_p = false;
+
+#pragma omp begin declare variant match(implementation={vendor(score(9):nvidia,amd)},user={condition(score(17):enable_p)})
+#pragma omp begin declare variant match(implementation={vendor(score(13):"nvidia")},user={condition(enable2_p)})
+int quux (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-3.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-3.C
new file mode 100644
index 00000000000..2f5a23f5594
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-3.C
@@ -0,0 +1,6 @@
+#pragma omp begin declare variant match(implementation={vendor(amd,nvidia),atomic_default_mem_order(relaxed)})
+#pragma omp begin declare variant match(implementation={atomic_default_mem_order(acq_rel)})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-4.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-4.C
new file mode 100644
index 00000000000..5e1478132e6
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-4.C
@@ -0,0 +1,10 @@
+#pragma omp begin declare variant match(implementation={vendor(amd,nvidia),atomic_default_mem_order(acq_rel)})
+#pragma omp begin declare variant match(implementation={atomic_default_mem_order(acq_rel)})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+
+int foo (int c) { return c; }
+
+int main() { foo (5); return 0; }
diff --git a/gcc/testsuite/g++.dg/gomp/delim-declare-variant-5.C b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-5.C
new file mode 100644
index 00000000000..25b08697dc2
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/delim-declare-variant-5.C
@@ -0,0 +1,4 @@
+#pragma omp begin declare variant match(device={arch("super power")})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+
diff --git a/gcc/tree-cfg.cc b/gcc/tree-cfg.cc
index 1ab18fa6b0f..48492951931 100644
--- a/gcc/tree-cfg.cc
+++ b/gcc/tree-cfg.cc
@@ -8300,6 +8300,10 @@ dump_function_to_file (tree fndecl, FILE *file, dump_flags_t flags)
 		  print_omp_context_selector (file, TREE_VALUE (a),
 					      dump_flags);
 		}
+	      else if (!strcmp (IDENTIFIER_POINTER (name),
+				"omp declare variant overload"))
+		print_omp_context_selector (file, TREE_VALUE (chain),
+					    dump_flags);
 	      else
 		print_generic_expr (file, TREE_VALUE (chain), dump_flags);
 	      fprintf (file, ")");
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 65e51b939a2..0ffdbb20e01 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1994,8 +1994,9 @@ struct GTY(()) tree_function_decl {
   unsigned has_debug_args_flag : 1;
   unsigned versioned_function : 1;
   unsigned replaceable_operator : 1;
+  unsigned omp_variant : 1;
 
-  /* 11 bits left for future expansion.  */
+  /* 10 bits left for future expansion.  */
   /* 32 bits on 64-bit HW.  */
 };
 
diff --git a/gcc/tree.h b/gcc/tree.h
index 086b55f0375..019deecca40 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -3505,6 +3505,13 @@ extern vec<tree, va_gc> **decl_debug_args_insert (tree);
 #define DECL_FUNCTION_VERSIONED(NODE)\
    (FUNCTION_DECL_CHECK (NODE)->function_decl.versioned_function)
 
+/* In FUNCTION_DECL, a OpenMP 'declare variant' function from a
+   'begin/end declare variant' block.  These have the same name as some base
+   function, so this flag is used to disambiguate them (similar to
+   DECL_FUNCTION_VERSIONED).  */
+#define DECL_FUNCTION_OMP_VARIANT(NODE) \
+   (FUNCTION_DECL_CHECK (NODE)->function_decl.omp_variant)
+
 /* In FUNCTION_DECL, this is set if this function is a C++ constructor.
    Devirtualization machinery uses this knowledge for determing type of the
    object constructed.  Also we assume that constructor address is not
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-10.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-10.C
new file mode 100644
index 00000000000..2bc4fbc5260
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-10.C
@@ -0,0 +1,19 @@
+#pragma omp begin declare variant match(implementation={dynamic_allocators,unified_shared_memory})
+#pragma omp begin declare variant match(implementation={reverse_offload})
+int foo (int c)
+{
+  return c+5;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int foo (int c)
+{
+  return c+3;
+}
+
+int main(void)
+{
+  foo (6);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-13.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-13.C
new file mode 100644
index 00000000000..bbdf7fa6811
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-13.C
@@ -0,0 +1,51 @@
+#include <stdio.h>
+
+int foo(int c)
+{
+  return c+1;
+}
+
+int foo(int c, int d)
+{
+  return c+d;
+}
+
+int foo(int c, int d, int e, int f = 10)
+{
+  return c+d+e+f;
+}
+
+#pragma omp begin declare variant match(construct={target})
+int foo(int c)
+{
+  return (c+1)*2;
+}
+
+int foo(int c, int d)
+{
+  return (c+d)*2;
+}
+
+int foo(int c, int d, int e, int f = 10)
+{
+  return (c+d+e+f)*2;
+}
+#pragma omp end declare variant
+
+int main()
+{
+  printf ("foo(5) = %d\n", foo(5));
+  printf ("foo(5,2) = %d\n", foo(5,2));
+  printf ("foo(5,2,3) = %d\n", foo(5,2,3));
+  printf ("foo(5,2,3,4) = %d\n", foo(5,2,3,4));
+
+#pragma omp target
+  {
+    printf ("foo(5) = %d\n", foo(5));
+    printf ("foo(5,2) = %d\n", foo(5,2));
+    printf ("foo(5,2,3) = %d\n", foo(5,2,3));
+    printf ("foo(5,2,3,4) = %d\n", foo(5,2,3,4));
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-14.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-14.C
new file mode 100644
index 00000000000..99739d8cae1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-14.C
@@ -0,0 +1,57 @@
+#include <stdio.h>
+
+template<typename T>
+T foo(T c)
+{
+  return c+1;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+  return c+d;
+}
+
+template<typename T, int U>
+T foo(T c, T d, T e, T f = U)
+{
+  return c+d+e+f;
+}
+
+#pragma omp begin declare variant match(construct={target})
+template<typename T>
+T foo(T c)
+{
+  return (c+1)*2;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+  return (c+d)*2;
+}
+
+template<typename T, int U>
+T foo(T c, T d, T e, T f = U)
+{
+  return (c+d+e+f)*2;
+}
+#pragma omp end declare variant
+
+int main()
+{
+  printf ("foo(5) = %d\n", foo<int>(5));
+  printf ("foo(5,2) = %d\n", foo<int>(5,2));
+  printf ("foo(5,2,3) = %d\n", foo<int, 5>(5,2,3));
+  printf ("foo(5,2,3,4) = %d\n", foo<int, 10>(5,2,3,4));
+
+#pragma omp target
+  {
+    printf ("foo(5) = %d\n", foo<int>(5));
+    printf ("foo(5,2) = %d\n", foo<int>(5,2));
+    printf ("foo(5,2,3) = %d\n", foo<int, 5>(5,2,3));
+    printf ("foo(5,2,3,4) = %d\n", foo<int, 10>(5,2,3,4));
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-15.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-15.C
new file mode 100644
index 00000000000..1f953e25ec0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-15.C
@@ -0,0 +1,55 @@
+#include <stdio.h>
+
+template<typename T>
+T foo(T c)
+{
+  return c+1;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+  return c+d;
+}
+
+template<typename T, typename S>
+T foo(T c, S d)
+{
+  return c+d+2;
+}
+
+#pragma omp begin declare variant match(construct={target})
+template<typename T>
+T foo(T c)
+{
+  return (c+1)*2;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+  return (c+d)*2;
+}
+
+template<typename T, typename S>
+T foo(T c, S d)
+{
+  return (c+d+2)*2;
+}
+#pragma omp end declare variant
+
+int main()
+{
+  printf ("foo<int>(5) = %d\n", foo<int>(5));
+  printf ("foo<int>(5,2) = %d\n", foo<int>(5,2));
+  printf ("foo<int,long>(5,2) = %d\n", foo<int, long>(5,2));
+
+#pragma omp target
+  {
+    printf ("foo<int>(5) = %d\n", foo<int>(5));
+    printf ("foo<int>(5,2) = %d\n", foo<int>(5,2));
+    printf ("foo<int,long>(5,2) = %d\n", foo<int, long>(5,2));
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-16.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-16.C
new file mode 100644
index 00000000000..22551d24469
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-16.C
@@ -0,0 +1,30 @@
+#include <cstdio>
+
+class foo {
+public:
+  int bar (int c, int d) { return c + d; }
+  static int qux (int c, int d) { return c + d + 5; }
+
+#if 1
+#pragma omp begin declare variant match(construct={target})
+  int bar (int c, int d) { return c + d + 3; }
+  static int qux (int c, int d) { return c + d + 7; }
+#pragma omp end declare variant
+#endif
+};
+
+int main()
+{
+  foo fv;
+
+  printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+  printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+
+#pragma omp target
+  {
+    printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+    printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-17.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-17.C
new file mode 100644
index 00000000000..0d70cf51d50
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-17.C
@@ -0,0 +1,29 @@
+#include <cstdio>
+
+class foo {
+public:
+  int t_bar (int c, int d) { return c + d; }
+  static int t_qux (int c, int d) { return c + d + 5; }
+
+#pragma omp declare variant (t_bar) match(construct={target})
+  int bar (int c, int d) { return c + d + 3; }
+
+#pragma omp declare variant (t_qux) match(construct={target})
+  static int qux (int c, int d) { return c + d + 7; }
+};
+
+int main()
+{
+  foo fv;
+
+  printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+  printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+
+#pragma omp target
+  {
+    printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+    printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-18.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-18.C
new file mode 100644
index 00000000000..1db28324bbe
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-18.C
@@ -0,0 +1,46 @@
+#include <cassert>
+#include <cstdio>
+
+#if 0
+#pragma omp begin declare variant match(device={arch(x86)})
+template<typename T>
+T foo (T v)
+{
+  return v + 3;
+}
+#pragma omp end declare variant
+#endif
+
+#pragma omp begin declare variant match(device={arch(powerpc)})
+template<typename T>
+T foo (T v)
+{
+  return v + 5;
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={arch(aarch64)})
+template<typename T>
+T foo (T v)
+{
+  return v + 7;
+}
+#pragma omp end declare variant
+
+template<typename T>
+T foo (T v)
+{
+  return v + 1;
+}
+
+int main ()
+{
+  int res;
+#pragma omp dispatch novariants(1)
+  {
+    res = foo<int> (3);
+  }
+  printf ("dispatch foo(3)=%d\n", res);
+  printf ("regular foo(3)=%d\n", foo<int> (3));
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-19.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-19.C
new file mode 100644
index 00000000000..077672299ff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-19.C
@@ -0,0 +1,41 @@
+#include <cassert>
+#include <cstdio>
+
+template<typename T>
+T foo_x86 (T v)
+{
+  return v + 3;
+}
+
+template<typename T>
+T foo_powerpc (T v)
+{
+  return v + 5;
+}
+
+template<typename T>
+T foo_aarch64 (T v)
+{
+  return v + 7;
+}
+
+#pragma omp declare variant (foo_x86) match(device={arch(x86)})
+#pragma omp declare variant (foo_powerpc) match(device={arch(powerpc)})
+#pragma omp declare variant (foo_aarch64) match(device={arch(arm)})
+template<typename T>
+T foo (T v)
+{
+  return v + 1;
+}
+
+int main ()
+{
+  int res;
+#pragma omp dispatch novariants(1)
+  {
+    res = foo<int> (3);
+  }
+  printf ("dispatch foo(3)=%d\n", res);
+  printf ("regular foo(3)=%d\n", foo<int> (3));
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-20.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-20.C
new file mode 100644
index 00000000000..b1d4f500d1c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-20.C
@@ -0,0 +1,39 @@
+#include <cassert>
+#include <cstdio>
+
+typedef int T;
+
+T foo_x86 (T v)
+{
+  return v + 3;
+}
+
+T foo_powerpc (T v)
+{
+  return v + 5;
+}
+
+T foo_aarch64 (T v)
+{
+  return v + 7;
+}
+
+#pragma omp declare variant (foo_x86) match(device={arch(x86)})
+#pragma omp declare variant (foo_powerpc) match(device={arch(powerpc)})
+#pragma omp declare variant (foo_aarch64) match(device={arch(aarch64)})
+T foo (T v)
+{
+  return v + 1;
+}
+
+int main ()
+{
+  int res;
+#pragma omp dispatch novariants(1)
+  {
+    res = foo (3);
+  }
+  printf ("dispatch foo(3)=%d\n", res);
+  printf ("regular foo(3)=%d\n", foo (3));
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-21.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-21.C
new file mode 100644
index 00000000000..7a26c9fdae9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-21.C
@@ -0,0 +1,56 @@
+#include <cstdio>
+
+template<typename T>
+class foo
+{
+  T var;
+public:
+  void set_var(T c) { var = c; }
+  T get_var () { return var + 1; }
+
+#pragma omp begin declare variant match(construct={target})
+  T get_var () { return var + 3; }
+#pragma omp end declare variant
+};
+
+void inty ()
+{
+  foo<int> fv;
+  fv.set_var (6);
+
+  printf ("fv.get_var () = %d\n", fv.get_var ());
+
+  int res;
+
+#pragma omp target map(from: res)
+  {
+    res = fv.get_var ();
+  }
+
+  printf ("fv.get_var () [target] = %d\n", res);
+}
+
+void longy ()
+{
+  foo<long> fvl;
+  fvl.set_var (7);
+
+  printf ("fvl.get_var () = %ld\n", fvl.get_var ());
+
+  long res;
+
+#pragma omp target map(from: res)
+  {
+    res = fvl.get_var ();
+  }
+
+  printf ("fvl.get_var () [target] = %ld\n", res);
+}
+
+int main ()
+{
+  inty ();
+  longy ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-6.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-6.C
new file mode 100644
index 00000000000..ab99f874f86
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-6.C
@@ -0,0 +1,37 @@
+#include <stdio.h>
+
+#pragma omp begin declare variant match(construct={target})
+int foo (void)
+{
+  return 7;
+}
+#pragma omp begin declare variant match(construct={teams})
+int foo (void)
+{
+  return 9;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int foo (void)
+{
+  return 5;
+}
+
+int main (void)
+{
+  int c = foo ();
+  printf ("host c=%d\n", c);
+
+#pragma omp target map(from:c)
+  {
+    c = foo ();
+  }
+  printf ("target c=%d\n", c);
+#pragma omp target teams map(from:c)
+  {
+    c = foo ();
+  }
+  printf ("target teams c=%d\n", c);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-7.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-7.C
new file mode 100644
index 00000000000..9dd697f9622
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-7.C
@@ -0,0 +1,38 @@
+#include <stdio.h>
+
+int p_foo (int c, int d)
+{
+  return (c+d) * 2;
+}
+
+int t_foo (int c, int d)
+{
+  return (c+d) * 3;
+}
+
+#pragma omp declare variant (p_foo) match(construct={parallel})
+#pragma omp declare variant (t_foo) match(construct={target})
+int foo (int c, int d)
+{
+  return c+d;
+}
+
+int main()
+{
+  printf ("host: %d\n", foo (3, 4));
+
+#pragma omp parallel
+  {
+#pragma omp single
+    {
+      printf ("parallel: %d\n", foo (3, 4));
+    }
+  }
+
+#pragma omp target
+  {
+    printf ("target: %d\n", foo (3, 4));
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-8.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-8.C
new file mode 100644
index 00000000000..b674474ddcc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-8.C
@@ -0,0 +1,61 @@
+#include <stdio.h>
+
+#define VARIANTS
+
+#ifdef VARIANTS
+#pragma omp begin declare variant match(construct={target})
+template<typename T>
+T foo (void)
+{
+  return 7;
+}
+#pragma omp begin declare variant match(construct={teams})
+template<typename T>
+T foo (void)
+{
+  return 9;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+#endif
+
+template<typename T>
+T foo (void)
+{
+  return 5;
+}
+
+int main (void)
+{
+  int c = foo<int> ();
+  printf ("host c=%d\n", c);
+
+#pragma omp target map(from:c)
+  {
+    c = foo<int> ();
+  }
+  printf ("target c=%d\n", c);
+#pragma omp target teams map(from:c)
+  {
+    c = foo<int> ();
+  }
+  printf ("target teams c=%d\n", c);
+
+  long cl = foo<long> ();
+  printf ("host cl=%ld\n", cl);
+
+#pragma omp target map(from:cl)
+  {
+    cl = foo<long> ();
+  }
+  printf ("target cl=%ld\n", cl);
+#pragma omp target teams map(from:cl)
+  {
+    cl = foo<long> ();
+  }
+  printf ("target teams cl=%ld\n", cl);
+
+ // int q = foo ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/delim-declare-variant-9.C b/libgomp/testsuite/libgomp.c++/delim-declare-variant-9.C
new file mode 100644
index 00000000000..135660f559b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/delim-declare-variant-9.C
@@ -0,0 +1,54 @@
+#include <stdio.h>
+
+template<typename T>
+T foo_targ (void)
+{
+  return 7;
+}
+
+template<typename T>
+T foo_targteams (void)
+{
+  return 9;
+}
+
+#pragma omp declare variant(foo_targ<T>) match(construct={target})
+#pragma omp declare variant(foo_targteams<T>) match(construct={target,teams})
+template<typename T>
+T foo (void)
+{
+  return 5;
+}
+
+int main (void)
+{
+  int c = foo<int> ();
+  printf ("host c=%d\n", c);
+
+#pragma omp target map(from:c)
+  {
+    c = foo<int> ();
+  }
+  printf ("target c=%d\n", c);
+#pragma omp target teams map(from:c)
+  {
+    c = foo<int> ();
+  }
+  printf ("target teams c=%d\n", c);
+
+  long cl = foo<long> ();
+  printf ("host cl=%ld\n", cl);
+
+#pragma omp target map(from:cl)
+  {
+    cl = foo<long> ();
+  }
+  printf ("target cl=%ld\n", cl);
+#pragma omp target teams map(from:cl)
+  {
+    cl = foo<long> ();
+  }
+  printf ("target teams cl=%ld\n", cl);
+
+  return 0;
+}
