diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc
index 4dd2eecbea5..fab9b5b8b23 100644
--- a/gcc/c-family/c-attribs.cc
+++ b/gcc/c-family/c-attribs.cc
@@ -571,6 +571,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 adjust_args need_device_ptr", 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,
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 620a3c1353a..5a0ed636677 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4300,8 +4300,8 @@ const struct c_omp_directive c_omp_directives[] = {
     C_OMP_DIR_DECLARATIVE, false },
   { "depobj", nullptr, nullptr, PRAGMA_OMP_DEPOBJ,
     C_OMP_DIR_STANDALONE, false },
-  /* { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH,
-    C_OMP_DIR_CONSTRUCT, false },  */
+  { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH,
+    C_OMP_DIR_DECLARATIVE, false },
   { "distribute", nullptr, nullptr, PRAGMA_OMP_DISTRIBUTE,
     C_OMP_DIR_CONSTRUCT, true },
   { "end", "assumes", nullptr, PRAGMA_OMP_END,
diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc
index ed2a7a00e9e..040370cbb6f 100644
--- a/gcc/c-family/c-pragma.cc
+++ b/gcc/c-family/c-pragma.cc
@@ -1526,6 +1526,7 @@ static const struct omp_pragma_def omp_pragmas[] = {
   { "cancellation", PRAGMA_OMP_CANCELLATION_POINT },
   { "critical", PRAGMA_OMP_CRITICAL },
   { "depobj", PRAGMA_OMP_DEPOBJ },
+  { "dispatch", PRAGMA_OMP_DISPATCH },
   { "error", PRAGMA_OMP_ERROR },
   { "end", PRAGMA_OMP_END },
   { "flush", PRAGMA_OMP_FLUSH },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 2ebde06c471..6b6826b2426 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -55,6 +55,7 @@ enum pragma_kind {
   PRAGMA_OMP_CRITICAL,
   PRAGMA_OMP_DECLARE,
   PRAGMA_OMP_DEPOBJ,
+  PRAGMA_OMP_DISPATCH,
   PRAGMA_OMP_DISTRIBUTE,
   PRAGMA_OMP_ERROR,
   PRAGMA_OMP_END,
@@ -135,9 +136,11 @@ enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_LINK,
   PRAGMA_OMP_CLAUSE_MAP,
   PRAGMA_OMP_CLAUSE_MERGEABLE,
+  PRAGMA_OMP_CLAUSE_NOCONTEXT,
   PRAGMA_OMP_CLAUSE_NOGROUP,
   PRAGMA_OMP_CLAUSE_NONTEMPORAL,
   PRAGMA_OMP_CLAUSE_NOTINBRANCH,
+  PRAGMA_OMP_CLAUSE_NOVARIANTS,
   PRAGMA_OMP_CLAUSE_NOWAIT,
   PRAGMA_OMP_CLAUSE_NUM_TASKS,
   PRAGMA_OMP_CLAUSE_NUM_TEAMS,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index a681438cbbe..61403093382 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -1747,6 +1747,8 @@ static void c_parser_omp_assumption_clauses (c_parser *, bool);
 static void c_parser_omp_allocate (c_parser *);
 static void c_parser_omp_assumes (c_parser *);
 static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *);
+static tree
+c_parser_omp_dispatch (location_t, c_parser *);
 static void c_parser_oacc_routine (c_parser *, enum pragma_context);
 
 /* These Objective-C parser functions are only ever called when
@@ -15090,6 +15092,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	case 'n':
 	  if (!strcmp ("no_create", p))
 	    result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+	  else if (!strcmp ("nocontext", p))
+	    result = PRAGMA_OMP_CLAUSE_NOCONTEXT;
 	  else if (!strcmp ("nogroup", p))
 	    result = PRAGMA_OMP_CLAUSE_NOGROUP;
 	  else if (!strcmp ("nohost", p))
@@ -15098,6 +15102,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
 	  else if (!strcmp ("notinbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
+	  else if (!strcmp ("novariants", p))
+	    result = PRAGMA_OMP_CLAUSE_NOVARIANTS;
 	  else if (!strcmp ("nowait", p))
 	    result = PRAGMA_OMP_CLAUSE_NOWAIT;
 	  else if (!strcmp ("num_gangs", p))
@@ -19365,6 +19371,60 @@ c_parser_omp_clause_partial (c_parser *parser, tree list)
   return c;
 }
 
+/* OpenMP 5.1
+   novariants ( scalar-expression ) */
+
+static tree
+c_parser_omp_clause_novariants (c_parser *parser, tree list)
+{
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
+
+  location_t loc = c_parser_peek_token (parser)->location;
+  c_expr expr = c_parser_expr_no_commas (parser, NULL);
+  tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value;
+  t = c_objc_common_truthvalue_conversion (loc, t);
+  t = c_fully_fold (t, false, NULL);
+  parens.skip_until_found_close (parser);
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants");
+
+  tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS);
+  OMP_CLAUSE_NOVARIANTS_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  list = c;
+
+  return list;
+}
+
+/* OpenMP 5.1
+   nocontext ( scalar-expression ) */
+
+static tree
+c_parser_omp_clause_nocontext (c_parser *parser, tree list)
+{
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
+
+  location_t loc = c_parser_peek_token (parser)->location;
+  c_expr expr = c_parser_expr_no_commas (parser, NULL);
+  tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value;
+  t = c_objc_common_truthvalue_conversion (loc, t);
+  t = c_fully_fold (t, false, NULL);
+  parens.skip_until_found_close (parser);
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext");
+
+  tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT);
+  OMP_CLAUSE_NOCONTEXT_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  list = c;
+
+  return list;
+}
+
 /* OpenMP 5.0:
    detach ( event-handle ) */
 
@@ -19984,6 +20044,14 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  c_name = "partial";
 	  clauses = c_parser_omp_clause_partial (parser, clauses);
 	  break;
+	case PRAGMA_OMP_CLAUSE_NOVARIANTS:
+	  c_name = "novariants";
+	  clauses = c_parser_omp_clause_novariants (parser, clauses);
+	  break;
+	case PRAGMA_OMP_CLAUSE_NOCONTEXT:
+	  c_name = "nocontext";
+	  clauses = c_parser_omp_clause_nocontext (parser, clauses);
+	  break;
 	default:
 	  c_parser_error (parser, "expected an OpenMP clause");
 	  goto saw_error;
@@ -23794,6 +23862,202 @@ c_parser_omp_scope (location_t loc, c_parser *parser, bool *if_p)
   return add_stmt (stmt);
 }
 
+// Adapted from c_parser_expr_no_commas
+static tree
+c_parser_omp_dispatch_body (c_parser *parser)
+{
+  struct c_expr lhs, rhs, ret;
+  struct c_expr orig_expr;
+  location_t expr_loc = c_parser_peek_token (parser)->location;
+  source_range tok_range = c_parser_peek_token (parser)->get_range ();
+  location_t sizeof_arg_loc[3];
+  tree sizeof_arg[3];
+  vec<tree, va_gc> *exprlist;
+  vec<location_t> arg_loc = vNULL;
+  vec<tree, va_gc> *origtypes = NULL;
+  unsigned int literal_zero_mask;
+  location_t start;
+  location_t finish;
+
+  lhs = c_parser_conditional_expression (parser, NULL, NULL);
+  if (TREE_CODE (lhs.value) == CALL_EXPR)
+    return lhs.value;
+  else
+    {
+      location_t op_location = c_parser_peek_token (parser)->location;
+      if (!c_parser_require (parser, CPP_EQ, "expected %<=%>"))
+	return error_mark_node;
+
+      /* Parse function name*/
+      if (!c_parser_next_token_is (parser, CPP_NAME))
+	{
+	  c_parser_error (parser, "expected a function name");
+	  rhs.set_error ();
+	  return rhs.value;
+	}
+      expr_loc = c_parser_peek_token (parser)->location;
+      tree id = c_parser_peek_token (parser)->value;
+      c_parser_consume_token (parser);
+      if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+	return error_mark_node;
+
+      rhs.value = build_external_ref (expr_loc, id, true, &rhs.original_type);
+      set_c_expr_source_range (&rhs, tok_range);
+      /* Parse argument list  */
+      {
+	for (int i = 0; i < 3; i++)
+	  {
+	    sizeof_arg[i] = NULL_TREE;
+	    sizeof_arg_loc[i] = UNKNOWN_LOCATION;
+	  }
+	literal_zero_mask = 0;
+	if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+	  exprlist = NULL;
+	else
+	  exprlist = c_parser_expr_list (parser, true, false, &origtypes,
+					 sizeof_arg_loc, sizeof_arg, &arg_loc,
+					 &literal_zero_mask);
+	c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+      }
+      orig_expr = rhs;
+      mark_exp_read (rhs.value);
+      if (warn_sizeof_pointer_memaccess)
+	sizeof_pointer_memaccess_warning (sizeof_arg_loc, rhs.value, exprlist,
+					  sizeof_arg,
+					  sizeof_ptr_memacc_comptypes);
+      if (TREE_CODE (rhs.value) == FUNCTION_DECL)
+	{
+	  if (fndecl_built_in_p (rhs.value, BUILT_IN_MEMSET)
+	      && vec_safe_length (exprlist) == 3)
+	    {
+	      tree arg0 = (*exprlist)[0];
+	      tree arg2 = (*exprlist)[2];
+	      warn_for_memset (expr_loc, arg0, arg2, literal_zero_mask);
+	    }
+	  if (warn_absolute_value
+	      && fndecl_built_in_p (rhs.value, BUILT_IN_NORMAL)
+	      && vec_safe_length (exprlist) == 1)
+	    warn_for_abs (expr_loc, rhs.value, (*exprlist)[0]);
+	  if (parser->omp_for_parse_state
+	      && parser->omp_for_parse_state->in_intervening_code
+	      && omp_runtime_api_call (rhs.value))
+	    {
+	      error_at (expr_loc, "calls to the OpenMP runtime API are "
+				  "not permitted in intervening code");
+	      parser->omp_for_parse_state->fail = true;
+	    }
+	}
+
+      start = rhs.get_start ();
+      finish = parser->tokens_buf[0].get_finish ();
+      rhs.value = c_build_function_call_vec (expr_loc, arg_loc, rhs.value,
+					     exprlist, origtypes);
+      set_c_expr_source_range (&rhs, start, finish);
+      rhs.m_decimal = 0;
+
+      rhs.original_code = ERROR_MARK;
+      if (TREE_CODE (rhs.value) == INTEGER_CST
+	  && TREE_CODE (orig_expr.value) == FUNCTION_DECL
+	  && fndecl_built_in_p (orig_expr.value, BUILT_IN_CONSTANT_P))
+	rhs.original_code = C_MAYBE_CONST_EXPR;
+      rhs.original_type = NULL;
+      if (exprlist)
+	{
+	  release_tree_vector (exprlist);
+	  release_tree_vector (origtypes);
+	}
+      arg_loc.release ();
+
+      /* Build assignment */
+      rhs = convert_lvalue_to_rvalue (expr_loc, rhs, true, true);
+      ret.value
+	= build_modify_expr (op_location, lhs.value, lhs.original_type,
+			     NOP_EXPR, expr_loc, rhs.value, rhs.original_type);
+      ret.m_decimal = 0;
+      set_c_expr_source_range (&ret, lhs.get_start (), rhs.get_finish ());
+      ret.original_code = MODIFY_EXPR;
+      ret.original_type = NULL;
+      return ret.value;
+    }
+}
+
+/* OpenMP 5.1:
+   # pragma omp dispatch dispatch-clause[optseq] new-line
+     expression-stmt
+
+   LOC is the location of the #pragma.
+*/
+
+#define OMP_DISPATCH_CLAUSE_MASK                                               \
+  ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)                             \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND)                           \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS)                       \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT)                        \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)                    \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
+
+static tree
+c_parser_omp_dispatch (location_t loc, c_parser *parser)
+{
+  tree stmt = make_node (OMP_DISPATCH);
+  SET_EXPR_LOCATION (stmt, loc);
+  TREE_TYPE (stmt) = void_type_node;
+
+  OMP_DISPATCH_CLAUSES (stmt)
+    = c_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK,
+				"#pragma omp dispatch");
+
+  // Extract depend clauses and create taskwait
+  tree depend_clauses = NULL_TREE;
+  tree *depend_clauses_ptr = &depend_clauses;
+  for (tree c = OMP_DISPATCH_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)
+	{
+	  *depend_clauses_ptr = c;
+	  depend_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	}
+    }
+  if (depend_clauses != NULL_TREE)
+    {
+      tree stmt = make_node (OMP_TASK);
+      TREE_TYPE (stmt) = void_node;
+      OMP_TASK_CLAUSES (stmt) = depend_clauses;
+      OMP_TASK_BODY (stmt) = NULL_TREE;
+      SET_EXPR_LOCATION (stmt, loc);
+      add_stmt (stmt);
+    }
+
+  // Parse body as expression statement
+  loc = c_parser_peek_token (parser)->location;
+  tree dispatch_body = c_parser_omp_dispatch_body (parser);
+  if (dispatch_body == error_mark_node)
+    {
+      inform (loc, "%<#pragma omp dispatch%> must be followed by a function "
+		   "call with optional assignment");
+      c_parser_skip_to_end_of_block_or_statement (parser);
+      return NULL_TREE;
+    }
+
+  // Walk the tree to find the dispatch function call and wrap it into an IFN
+  gcc_assert (TREE_CODE (dispatch_body) == CALL_EXPR
+	      || TREE_CODE (dispatch_body) == MODIFY_EXPR);
+  tree *dispatch_call = TREE_CODE (dispatch_body) == MODIFY_EXPR
+			  ? &TREE_OPERAND (dispatch_body, 1)
+			  : &dispatch_body;
+  if (TREE_CODE (*dispatch_call) == FLOAT_EXPR
+      || TREE_CODE (*dispatch_call) == CONVERT_EXPR)
+    dispatch_call = &TREE_OPERAND (*dispatch_call, 0);
+  *dispatch_call = build_call_expr_internal_loc (
+    loc, IFN_GOMP_DISPATCH,
+    TREE_TYPE (TREE_TYPE (CALL_EXPR_FN (*dispatch_call))), 1, *dispatch_call);
+
+  c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>");
+  OMP_DISPATCH_BODY (stmt) = dispatch_body;
+
+  return add_stmt (stmt);
+}
+
 /* OpenMP 3.0:
    # pragma omp task task-clause[optseq] new-line
 
@@ -24774,6 +25038,10 @@ check_clauses:
 
    OpenMP 5.0:
    # pragma omp declare variant (identifier) match(context-selector) new-line
+
+   OpenMP 5.1
+   # pragma omp declare variant (identifier) match(context-selector) \
+      adjust_args(adjust-op:argument-list) new-line
    */
 
 #define OMP_DECLARE_SIMD_CLAUSE_MASK				\
@@ -25237,77 +25505,223 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
 
   parens.require_close (parser);
 
-  if (c_parser_next_token_is (parser, CPP_COMMA)
-      && c_parser_peek_2nd_token (parser)->type == CPP_NAME)
-    c_parser_consume_token (parser);
+  vec<tree> adjust_args_list = vNULL;
+  bool has_match = false, has_adjust_args = false;
+  location_t adjust_args_loc = UNKNOWN_LOCATION;
+  tree need_device_ptr_list = make_node (TREE_LIST);
 
-  const char *clause = "";
-  location_t match_loc = c_parser_peek_token (parser)->location;
-  if (c_parser_next_token_is (parser, CPP_NAME))
-    clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
-  if (strcmp (clause, "match"))
+  do
     {
-      c_parser_error (parser, "expected %<match%>");
-      goto fail;
-    }
+      if (c_parser_next_token_is (parser, CPP_COMMA)
+	  && c_parser_peek_2nd_token (parser)->type == CPP_NAME)
+	c_parser_consume_token (parser);
 
-  c_parser_consume_token (parser);
+      const char *clause = "";
+      location_t match_loc = c_parser_peek_token (parser)->location;
+      if (c_parser_next_token_is (parser, CPP_NAME))
+	clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
 
-  if (!parens.require_open (parser))
-    goto fail;
+      enum clause
+      {
+	match,
+	adjust_args
+      } ccode;
 
-  if (parms == NULL_TREE)
-    parms = error_mark_node;
-
-  tree ctx = c_parser_omp_context_selector_specification (parser, parms);
-  if (ctx == error_mark_node)
-    goto fail;
-  ctx = omp_check_context_selector (match_loc, ctx);
-  if (ctx != error_mark_node && variant != error_mark_node)
-    {
-      if (TREE_CODE (variant) != FUNCTION_DECL)
+      if (strcmp (clause, "match") == 0)
+	ccode = match;
+      else if (strcmp (clause, "adjust_args") == 0)
 	{
-	  error_at (token->location, "variant %qD is not a function", variant);
-	  variant = error_mark_node;
+	  ccode = adjust_args;
+	  adjust_args_loc = match_loc;
 	}
-      else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
-					  OMP_TRAIT_CONSTRUCT_SIMD)
-	       && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
+      else
 	{
-	  error_at (token->location, "variant %qD and base %qD have "
-				     "incompatible types", variant, fndecl);
-	  variant = error_mark_node;
+	  c_parser_error (parser, "expected %<match%> or %<adjust_args%>");
+	  goto fail;
 	}
-      else if (fndecl_built_in_p (variant)
-	       && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
-			    "__builtin_", strlen ("__builtin_")) == 0
-		   || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
-			       "__sync_", strlen ("__sync_")) == 0
-		   || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
-			       "__atomic_", strlen ("__atomic_")) == 0))
+
+      c_parser_consume_token (parser);
+
+      if (!parens.require_open (parser))
+	goto fail;
+
+      if (parms == NULL_TREE)
+	parms = error_mark_node;
+
+      if (ccode == match)
 	{
-	  error_at (token->location, "variant %qD is a built-in", variant);
-	  variant = error_mark_node;
-	}
-      if (variant != error_mark_node)
-	{
-	  C_DECL_USED (variant) = 1;
-	  tree construct
-	    = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
-	  omp_mark_declare_variant (match_loc, variant, construct);
-	  if (omp_context_selector_matches (ctx))
+	  has_match = true;
+	  tree ctx
+	    = c_parser_omp_context_selector_specification (parser, parms);
+	  if (ctx == error_mark_node)
+	    goto fail;
+	  ctx = omp_check_context_selector (match_loc, ctx);
+	  if (ctx != error_mark_node && variant != error_mark_node)
 	    {
-	      tree attr
-		= tree_cons (get_identifier ("omp declare variant base"),
-			     build_tree_list (variant, ctx),
-			     DECL_ATTRIBUTES (fndecl));
-	      DECL_ATTRIBUTES (fndecl) = attr;
+	      if (TREE_CODE (variant) != FUNCTION_DECL)
+		{
+		  error_at (token->location, "variant %qD is not a function",
+			    variant);
+		  variant = error_mark_node;
+		}
+	      else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+						  OMP_TRAIT_CONSTRUCT_SIMD)
+		       && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
+		{
+		  error_at (token->location,
+			    "variant %qD and base %qD have "
+			    "incompatible types",
+			    variant, fndecl);
+		  variant = error_mark_node;
+		}
+	      else if (fndecl_built_in_p (variant)
+		       && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+				    "__builtin_", strlen ("__builtin_"))
+			     == 0
+			   || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+				       "__sync_", strlen ("__sync_"))
+				== 0
+			   || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+				       "__atomic_", strlen ("__atomic_"))
+				== 0))
+		{
+		  error_at (token->location, "variant %qD is a built-in",
+			    variant);
+		  variant = error_mark_node;
+		}
+	      if (variant != error_mark_node)
+		{
+		  C_DECL_USED (variant) = 1;
+		  tree construct
+		    = omp_get_context_selector_list (ctx,
+						     OMP_TRAIT_SET_CONSTRUCT);
+		  omp_mark_declare_variant (match_loc, variant, construct);
+		  if (omp_context_selector_matches (ctx))
+		    {
+		      tree attr = tree_cons (get_identifier (
+					       "omp declare variant base"),
+					     build_tree_list (variant, ctx),
+					     DECL_ATTRIBUTES (fndecl));
+		      DECL_ATTRIBUTES (fndecl) = attr;
+		    }
+		}
 	    }
 	}
+      else if (ccode == adjust_args)
+	{
+	  has_adjust_args = true;
+	  if (c_parser_next_token_is (parser, CPP_NAME)
+	      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+	    {
+	      const char *p
+		= IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+	      if (strcmp (p, "need_device_ptr") == 0
+		  || strcmp (p, "nothing") == 0)
+		{
+		  c_parser_consume_token (parser); // need_device_ptr
+		  c_parser_consume_token (parser); // :
+
+		  location_t loc = c_parser_peek_token (parser)->location;
+		  tree list
+		    = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_ERROR,
+						  NULL_TREE);
+
+		  tree arg;
+		  for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c))
+		    {
+		      tree decl = TREE_PURPOSE (c);
+		      int idx;
+		      for (arg = parms, idx = 0; arg != NULL;
+			   arg = TREE_CHAIN (arg), idx++)
+			if (arg == decl)
+			  break;
+		      if (arg == NULL_TREE)
+			{
+			  error_at (loc, "%qD is not a function argument",
+				    decl);
+			  goto fail;
+			}
+		      if (adjust_args_list.contains (arg))
+			{
+			  error_at (loc, "%qD is specified more than once",
+				    decl);
+			  goto fail;
+			}
+		      if (strcmp (p, "need_device_ptr") == 0
+			  && TREE_CODE (TREE_TYPE (arg)) != POINTER_TYPE)
+			{
+			  error_at (loc, "%qD is not a C pointer", decl);
+			  goto fail;
+			}
+		      adjust_args_list.safe_push (arg);
+		      if (strcmp (p, "need_device_ptr") == 0)
+			{
+			  need_device_ptr_list = chainon (
+			    need_device_ptr_list,
+			    build_tree_list (
+			      NULL_TREE,
+			      build_int_cst (
+				integer_type_node,
+				idx))); // Store 0-based argument index,
+					// as in gimplify_call_expr
+			}
+		    }
+		}
+	      else
+		{
+		  error_at (c_parser_peek_token (parser)->location,
+			    "expected %<nothing%> or %<need_device_ptr%>");
+		  goto fail;
+		}
+	    }
+	  else
+	    {
+	      error_at (c_parser_peek_token (parser)->location,
+			"expected %<nothing%> or %<need_device_ptr%> "
+			"followed by %<:%>");
+	      goto fail;
+	    }
+	}
+
+      parens.require_close (parser);
+  } while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL));
+  c_parser_skip_to_pragma_eol (parser);
+
+  if (has_adjust_args)
+    {
+      if (!has_match)
+	{
+	  error_at (
+	    adjust_args_loc,
+	    "an %<adjust_args%> clause can only be specified if the "
+	    "%<dispatch%> selector of the construct selector set appears "
+	    "in the %<match%> clause");
+	}
+      else
+	{
+	  tree attr = lookup_attribute ("omp declare variant base",
+					DECL_ATTRIBUTES (fndecl));
+	  tree ctx = TREE_VALUE (TREE_VALUE (attr));
+	  if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+					 OMP_TRAIT_CONSTRUCT_DISPATCH))
+	    error_at (
+	      adjust_args_loc,
+	      "an %<adjust_args%> clause can only be specified if the "
+	      "%<dispatch%> selector of the construct selector set appears "
+	      "in the %<match%> clause");
+	}
     }
 
-  parens.require_close (parser);
-  c_parser_skip_to_pragma_eol (parser);
+  if (TREE_CHAIN (need_device_ptr_list) != NULL_TREE
+      && variant != error_mark_node)
+    {
+      tree variant_decl = tree_strip_nop_conversions (variant);
+      DECL_ATTRIBUTES (variant_decl)
+	= tree_cons (get_identifier ("omp declare variant variant adjust_args"),
+		     build_tree_list (need_device_ptr_list,
+				      NULL_TREE /*need_device_addr */),
+		     DECL_ATTRIBUTES (variant_decl));
+    }
 }
 
 /* Finalize #pragma omp declare simd or #pragma omp declare variant
@@ -26127,7 +26541,6 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
   types.release ();
 }
 
-
 /* OpenMP 4.0
    #pragma omp declare simd declare-simd-clauses[optseq] new-line
    #pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -26135,7 +26548,11 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
    #pragma omp declare target new-line
 
    OpenMP 5.0
-   #pragma omp declare variant (identifier) match (context-selector)  */
+   #pragma omp declare variant (identifier) match (context-selector)
+
+   OpenMP 5.1
+   #pragma omp declare variant (identifier) match (context-selector) \
+      adjust_args(adjust-op:argument-list)  */
 
 static bool
 c_parser_omp_declare (c_parser *parser, enum pragma_context context)
@@ -27053,6 +27470,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
     case PRAGMA_OMP_UNROLL:
       stmt = c_parser_omp_unroll (loc, parser, if_p);
       break;
+    case PRAGMA_OMP_DISPATCH:
+      stmt = c_parser_omp_dispatch (loc, parser);
+      break;
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index ba6d96d26b2..fa4777dd35a 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -16333,6 +16333,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_NOHOST:
 	case OMP_CLAUSE_INDIRECT:
+	case OMP_CLAUSE_NOVARIANTS:
+	case OMP_CLAUSE_NOCONTEXT:
 	  pc = &OMP_CLAUSE_CHAIN (c);
 	  continue;
 
diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c
new file mode 100644
index 00000000000..393a44de8e0
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c
@@ -0,0 +1,32 @@
+/* Test parsing of OMP clause adjust_args */
+/* { dg-do compile } */
+
+int b;
+
+int f0 (void *a);
+int g (void *a);
+int f1 (int);
+
+#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f2 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */
+int f3 (int a);
+#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f4 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f5 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f6 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected identifier before '\\)' token" } */
+int f7 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' undeclared here \\(not in a function\\)" } */
+int f8 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */
+int f9 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */
+int f10 (int a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */
+int f11 (void *a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */
+int f12 (void *a);
+
diff --git a/gcc/testsuite/gcc.dg/gomp/dispatch-1.c b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c
new file mode 100644
index 00000000000..c8f45c12be6
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c
@@ -0,0 +1,53 @@
+/* Test parsing of #pragma omp dispatch */
+/* { dg-do compile } */
+
+int f0 (int);
+
+void f1 (void)
+{
+  int a, b;
+  double x;
+  struct {int a; int b;} s;
+  int arr[1];
+
+#pragma omp dispatch
+  int c = f0 (a);	/* { dg-error "expected expression before 'int'" } */
+#pragma omp dispatch
+  int f2 (int d);	/* { dg-error "expected expression before 'int'" } */
+#pragma omp dispatch
+  a = b;	/* { dg-error "expected '\\(' before ';' token" } */
+#pragma omp dispatch
+  s.a = f0(a) + b;	/* { dg-error "expected ';' before '\\+' token" } */
+#pragma omp dispatch
+  b = !f0(a);	/* { dg-error "expected a function name before '!' token" } */
+#pragma omp dispatch
+  s.b += f0(s.a);	/* { dg-error "expected '=' before '\\+=' token" } */
+#pragma omp dispatch
+#pragma omp threadprivate(a)	/* { dg-error "expected expression before '#pragma'" } */
+  a = f0(b);
+  
+#pragma omp dispatch nocontext(s) /* { dg-error "used struct type value where scalar is required" } */
+  f0(a);
+#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */
+  f0(a);
+#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */
+  f0(a);
+#pragma omp dispatch novariants(s) /* { dg-error "used struct type value where scalar is required" } */
+  f0(a);
+#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */
+  f0(a);
+#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */
+  f0(a);
+#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */
+  f0(a);
+#pragma omp dispatch device(x) /* { dg-error "expected integer expression before end of line" } */
+  f0(a);
+#pragma omp dispatch device(arr) /* { dg-error "expected integer expression before end of line" } */
+  f0(a);
+#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer nor an array" } */
+  f0(a);
+#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected identifier before '&' token" } */
+  f0(a);
+#pragma omp dispatch depend(inout: f0) /* { dg-error "'f0' is not lvalue expression nor array section in 'depend' clause" } */
+  f0(a);
+}
diff --git a/libgomp/testsuite/libgomp.c/dispatch-1.c b/libgomp/testsuite/libgomp.c/dispatch-1.c
new file mode 100644
index 00000000000..0efc075a859
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/dispatch-1.c
@@ -0,0 +1,76 @@
+// Adapted from OpenMP examples
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int baz (double *d_bv, const double *d_av, int n)
+{
+#pragma omp distribute parallel for
+  for (int i = 0; i < n; i++)
+    d_bv[i] = d_av[i] * i;
+  return -3;
+}
+
+int bar (double *d_bv, const double *d_av, int n)
+{
+#pragma omp target is_device_ptr(d_bv, d_av)
+  for (int i = 0; i < n; i++)
+    d_bv[i] = d_av[i] * i;
+  return -2;
+}
+
+#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av)
+#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
+int foo (double *bv, const double *av, int n)
+{
+  for (int i = 0; i < n; i++)
+    bv[i] = av[i] * i;
+  return -1;
+}
+
+int test (int n)
+{
+  const double e = 2.71828;
+
+  double *av = (double *) malloc (n * sizeof (*av));
+  double *bv = (double *) malloc (n * sizeof (*bv));
+  double *d_bv = (double *) malloc (n * sizeof (*d_bv));
+
+  for (int i = 0; i < n; i++)
+    {
+      av[i] = e * i;
+      bv[i] = 0.0;
+      d_bv[i] = 0.0;
+    }
+
+  int f, last_dev = omp_get_num_devices () - 1;
+#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024)
+  {
+    #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev)
+    f = foo (d_bv, av, n);
+  }
+
+  foo (bv, av, n);
+  for (int i = 0; i < n; i++)
+    {
+      if (d_bv[i] != bv[i])
+	{
+	  fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]);
+	  return 1;
+	}
+    }
+    return f;
+}
+
+int
+main (void)
+{
+  int ret = test(1023);
+  if (ret != -1) return 1;
+  ret = test(1024);
+  if (ret != -2) return 1;
+  ret = test(1025);
+  if (ret != -3) return 1;
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/dispatch-2.c b/libgomp/testsuite/libgomp.c/dispatch-2.c
new file mode 100644
index 00000000000..faa0d8a1d1c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/dispatch-2.c
@@ -0,0 +1,84 @@
+// Adapted from OpenMP examples
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int baz (double *d_bv, const double *d_av, int n);
+int bar (double *d_bv, const double *d_av, int n); 
+
+#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: f_bv, f_av)
+#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
+int foo (double *f_bv, const double *f_av, int n);
+
+int baz (double *bv, const double *av, int n);
+int bar (double *bv, const double *av, int n); 
+
+int foo (double *bv, const double *av, int n)
+{
+  for (int i = 0; i < n; i++)
+    bv[i] = av[i] * i;
+  return -1;
+}
+
+int baz (double *d_bv, const double *d_av, int n)
+{
+#pragma omp distribute parallel for
+  for (int i = 0; i < n; i++)
+    d_bv[i] = d_av[i] * i;
+  return -3;
+}
+
+int bar (double *d_bv, const double *d_av, int n)
+{
+#pragma omp target is_device_ptr(d_bv, d_av)
+  for (int i = 0; i < n; i++)
+    d_bv[i] = d_av[i] * i;
+  return -2;
+}
+
+int test (int n)
+{
+  const double e = 2.71828;
+
+  double *av = (double *) malloc (n * sizeof (*av));
+  double *bv = (double *) malloc (n * sizeof (*bv));
+  double *d_bv = (double *) malloc (n * sizeof (*d_bv));
+
+  for (int i = 0; i < n; i++)
+    {
+      av[i] = e * i;
+      bv[i] = 0.0;
+      d_bv[i] = 0.0;
+    }
+
+  int f, last_dev = omp_get_num_devices () - 1;
+#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024)
+  {
+    #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev)
+    f = foo (d_bv, av, n);
+  }
+
+  foo (bv, av, n);
+  for (int i = 0; i < n; i++)
+    {
+      if (d_bv[i] != bv[i])
+	{
+	  fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]);
+	  return 1;
+	}
+    }
+    return f;
+}
+
+int
+main (void)
+{
+  int ret = test(1023);
+  if (ret != -1) return 1;
+  ret = test(1024);
+  if (ret != -2) return 1;
+  ret = test(1025);
+  if (ret != -3) return 1;
+  return 0;
+}
