[PATCHv2] openmp: Add support for 'present' modifier

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 43427886ad4..057a2c03937 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14014,6 +14014,8 @@  c_parser_omp_variable_list (c_parser *parser,
 	      tree u = build_omp_clause (clause_loc, kind);
 	      OMP_CLAUSE_DECL (u) = t;
 	      OMP_CLAUSE_CHAIN (u) = list;
+	      if (kind == OMP_CLAUSE_FROM || kind == OMP_CLAUSE_TO)
+		OMP_CLAUSE_SET_MOTION_MODIFIER (u, OMP_CLAUSE_MOTION_NONE);
 	      list = u;
 	    }
 	}
@@ -14041,7 +14043,8 @@  c_parser_omp_variable_list (c_parser *parser,
 
 static tree
 c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
-			      tree list, bool allow_deref = false)
+			      tree list, bool allow_deref = false,
+			      bool allow_present = false)
 {
   /* The clauses location.  */
   location_t loc = c_parser_peek_token (parser)->location;
@@ -14049,7 +14052,26 @@  c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
   matching_parens parens;
   if (parens.require_open (parser))
     {
+      bool present = false;
+
+      if (allow_present)
+	{
+	   c_token *token = c_parser_peek_token (parser);
+
+	   if (token->type == CPP_NAME
+	       && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
+	       && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+	    {
+	      present = true;
+	      c_parser_consume_token (parser);
+	      c_parser_consume_token (parser);
+	    }
+	}
       list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
+
+      if (present)
+	for (tree clause = list; clause; clause = OMP_CLAUSE_CHAIN (clause))
+	  OMP_CLAUSE_SET_MOTION_MODIFIER (clause, OMP_CLAUSE_MOTION_PRESENT);
       parens.skip_until_found_close (parser);
     }
   return list;
@@ -14933,6 +14955,13 @@  c_parser_omp_clause_defaultmap (c_parser *parser, tree list)
 	goto invalid_behavior;
       break;
 
+    case 'p':
+      if (strcmp ("present", p) == 0)
+	behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+      else
+	goto invalid_behavior;
+      break;
+
     case 't':
       if (strcmp ("tofrom", p) == 0)
 	behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -17103,6 +17132,7 @@  c_parser_omp_clause_map (c_parser *parser, tree list)
 
   int always_modifier = 0;
   int close_modifier = 0;
+  int present_modifier = 0;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       c_token *tok = c_parser_peek_token (parser);
@@ -17134,11 +17164,21 @@  c_parser_omp_clause_map (c_parser *parser, tree list)
 	    }
 	  close_modifier++;
 	}
+      else if (strcmp ("present", p) == 0)
+	{
+	  if (present_modifier)
+	    {
+	      c_parser_error (parser, "too many %<present%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+	  present_modifier++;
+	}
       else
 	{
 	  c_parser_error (parser, "%<#pragma omp target%> with "
-				  "modifier other than %<always%> or "
-				  "%<close%> on %<map%> clause");
+				  "modifier other than %<always%>, %<close%> "
+				  "or %<present%> on %<map%> clause");
 	  parens.skip_until_found_close (parser);
 	  return list;
 	}
@@ -17150,14 +17190,25 @@  c_parser_omp_clause_map (c_parser *parser, tree list)
       && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
     {
       const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      int always_present_modifier = always_modifier && present_modifier;
+
       if (strcmp ("alloc", p) == 0)
-	kind = GOMP_MAP_ALLOC;
+	kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+	       : present_modifier ? GOMP_MAP_PRESENT_TO
+	       : always_modifier ? GOMP_MAP_ALWAYS_TO
+	       : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+	       : present_modifier ? GOMP_MAP_PRESENT_FROM
+	       : always_modifier ? GOMP_MAP_ALWAYS_FROM
+	       : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+	       : present_modifier ? GOMP_MAP_PRESENT_TOFROM
+	       : always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+	       : GOMP_MAP_TOFROM;
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else if (strcmp ("delete", p) == 0)
@@ -17418,7 +17469,8 @@  c_parser_omp_clause_device_type (c_parser *parser, tree list)
 static tree
 c_parser_omp_clause_to (c_parser *parser, tree list)
 {
-  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true);
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true,
+				       true);
 }
 
 /* OpenMP 4.0:
@@ -17427,7 +17479,8 @@  c_parser_omp_clause_to (c_parser *parser, tree list)
 static tree
 c_parser_omp_clause_from (c_parser *parser, tree list)
 {
-  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true);
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true,
+				       true);
 }
 
 /* OpenMP 4.0:
@@ -21744,11 +21797,18 @@  c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	    map_seen = 3;
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -21894,7 +21954,10 @@  c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	    map_seen = 3;
 	    break;
 	  case GOMP_MAP_TOFROM:
@@ -21905,6 +21968,14 @@  c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+	    map_seen = 3;
+	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
@@ -21992,6 +22063,8 @@  c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	  {
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_RELEASE:
 	  case GOMP_MAP_DELETE:
 	    map_seen = 3;
@@ -22004,6 +22077,14 @@  c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+	    map_seen = 3;
+	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
@@ -22249,11 +22330,18 @@  check_clauses:
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 1a124f5395e..720827a942a 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37644,11 +37644,33 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 
 static tree
 cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
-			bool allow_deref = false)
+			bool allow_deref = false, bool allow_present = false)
 {
   if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
-    return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
-					   allow_deref);
+    {
+      bool present = false;
+
+      if (allow_present)
+	{
+	   cp_token *token = cp_lexer_peek_token (parser->lexer);
+
+	   if (token->type == CPP_NAME
+	       && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
+	       && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+	    {
+	      present = true;
+	      cp_lexer_consume_token (parser->lexer);
+	      cp_lexer_consume_token (parser->lexer);
+	    }
+	}
+
+      list = cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
+					     allow_deref);
+
+      if (present)
+	for (tree clause = list; clause; clause = OMP_CLAUSE_CHAIN (clause))
+	  OMP_CLAUSE_SET_MOTION_MODIFIER (clause, OMP_CLAUSE_MOTION_PRESENT);
+    }
   return list;
 }
 
@@ -38700,6 +38722,13 @@  cp_parser_omp_clause_defaultmap (cp_parser *parser, tree list,
 	goto invalid_behavior;
       break;
 
+    case 'p':
+      if (strcmp ("present", p) == 0)
+	behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+      else
+	goto invalid_behavior;
+      break;
+
     case 't':
       if (strcmp ("tofrom", p) == 0)
 	behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -40453,6 +40482,7 @@  cp_parser_omp_clause_map (cp_parser *parser, tree list)
 
   bool always_modifier = false;
   bool close_modifier = false;
+  bool present_modifier = false;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -40489,11 +40519,24 @@  cp_parser_omp_clause_map (cp_parser *parser, tree list)
 	    }
 	  close_modifier = true;
 	}
+      else if (strcmp ("present", p) == 0)
+	{
+	  if (present_modifier)
+	    {
+	      cp_parser_error (parser, "too many %<present%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+	  present_modifier = true;
+       }
       else
 	{
 	  cp_parser_error (parser, "%<#pragma omp target%> with "
-				   "modifier other than %<always%> or "
-				   "%<close%> on %<map%> clause");
+				   "modifier other than %<always%>, %<close%> "
+				   "or %<present%> on %<map%> clause");
 	  cp_parser_skip_to_closing_parenthesis (parser,
 						 /*recovering=*/true,
 						 /*or_comma=*/false,
@@ -40509,15 +40552,25 @@  cp_parser_omp_clause_map (cp_parser *parser, tree list)
     {
       tree id = cp_lexer_peek_token (parser->lexer)->u.value;
       const char *p = IDENTIFIER_POINTER (id);
+      int always_present_modifier = always_modifier && present_modifier;
 
       if (strcmp ("alloc", p) == 0)
-	kind = GOMP_MAP_ALLOC;
+	kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+	       : present_modifier ? GOMP_MAP_PRESENT_TO
+	       : always_modifier ? GOMP_MAP_ALWAYS_TO
+	       : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+	       : present_modifier ? GOMP_MAP_PRESENT_FROM
+	       : always_modifier ? GOMP_MAP_ALWAYS_FROM
+	       : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+	       : present_modifier ? GOMP_MAP_PRESENT_TOFROM
+	       : always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+	       : GOMP_MAP_TOFROM;
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else
@@ -41295,12 +41348,12 @@  cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	    }
 	  else
 	    clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
-					      true);
+					      true, true);
 	  c_name = "to";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FROM:
 	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
-					    true);
+					    true, true);
 	  c_name = "from";
 	  break;
 	case PRAGMA_OMP_CLAUSE_UNIFORM:
@@ -45131,11 +45184,18 @@  cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	    map_seen = 3;
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -45238,7 +45298,10 @@  cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	    map_seen = 3;
 	    break;
 	  case GOMP_MAP_TOFROM:
@@ -45249,6 +45312,14 @@  cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
 	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+	    map_seen = 3;
+	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
@@ -45341,6 +45412,8 @@  cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
 	  {
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_RELEASE:
 	  case GOMP_MAP_DELETE:
 	    map_seen = 3;
@@ -45353,6 +45426,14 @@  cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
 	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+	    map_seen = 3;
+	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index c2df0b69b30..6ba76ebeea6 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -9993,11 +9993,18 @@  finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 164710fe98a..b0b439ff0f3 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1453,9 +1453,20 @@  show_omp_namelist (int list_type, gfc_omp_namelist *n)
 	  case OMP_MAP_TO: fputs ("to:", dumpfile); break;
 	  case OMP_MAP_FROM: fputs ("from:", dumpfile); break;
 	  case OMP_MAP_TOFROM: fputs ("tofrom:", dumpfile); break;
+	  case OMP_MAP_PRESENT_ALLOC: fputs ("present,alloc:", dumpfile); break;
+	  case OMP_MAP_PRESENT_TO: fputs ("present,to:", dumpfile); break;
+	  case OMP_MAP_PRESENT_FROM: fputs ("present,from:", dumpfile); break;
+	  case OMP_MAP_PRESENT_TOFROM:
+	    fputs ("present,tofrom:", dumpfile); break;
 	  case OMP_MAP_ALWAYS_TO: fputs ("always,to:", dumpfile); break;
 	  case OMP_MAP_ALWAYS_FROM: fputs ("always,from:", dumpfile); break;
 	  case OMP_MAP_ALWAYS_TOFROM: fputs ("always,tofrom:", dumpfile); break;
+	  case OMP_MAP_ALWAYS_PRESENT_TO:
+	    fputs ("always,present,to:", dumpfile); break;
+	  case OMP_MAP_ALWAYS_PRESENT_FROM:
+	    fputs ("always,present,from:", dumpfile); break;
+	  case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    fputs ("always,present,tofrom:", dumpfile); break;
 	  case OMP_MAP_DELETE: fputs ("delete:", dumpfile); break;
 	  case OMP_MAP_RELEASE: fputs ("release:", dumpfile); break;
 	  default: break;
@@ -1778,6 +1789,10 @@  show_omp_clauses (gfc_omp_clauses *omp_clauses)
 	  fputs ("inscan, ", dumpfile);
 	if (list_type == OMP_LIST_REDUCTION_TASK)
 	  fputs ("task, ", dumpfile);
+	if ((list_type == OMP_LIST_TO || list_type == OMP_LIST_FROM)
+	    && omp_clauses->lists[list_type]->u.motion_modifier
+	       == OMP_MOTION_PRESENT)
+	  fputs ("present:", dumpfile);
 	show_omp_namelist (list_type, omp_clauses->lists[list_type]);
 	fputc (')', dumpfile);
       }
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index a893ee06f3d..e4681f0443f 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1306,7 +1306,14 @@  enum gfc_omp_map_op
   OMP_MAP_RELEASE,
   OMP_MAP_ALWAYS_TO,
   OMP_MAP_ALWAYS_FROM,
-  OMP_MAP_ALWAYS_TOFROM
+  OMP_MAP_ALWAYS_TOFROM,
+  OMP_MAP_PRESENT_ALLOC,
+  OMP_MAP_PRESENT_TO,
+  OMP_MAP_PRESENT_FROM,
+  OMP_MAP_PRESENT_TOFROM,
+  OMP_MAP_ALWAYS_PRESENT_TO,
+  OMP_MAP_ALWAYS_PRESENT_FROM,
+  OMP_MAP_ALWAYS_PRESENT_TOFROM
 };
 
 enum gfc_omp_defaultmap
@@ -1340,6 +1347,12 @@  enum gfc_omp_linear_op
   OMP_LINEAR_UVAL
 };
 
+enum gfc_omp_motion_modifier
+{
+  OMP_MOTION_NONE,
+  OMP_MOTION_PRESENT
+};
+
 /* For use in OpenMP clauses in case we need extra information
    (aligned clause alignment, linear clause step, etc.).  */
 
@@ -1358,6 +1371,7 @@  typedef struct gfc_omp_namelist
 	  ENUM_BITFIELD (gfc_omp_linear_op) op:4;
 	  bool old_modifier;
 	} linear;
+      gfc_omp_motion_modifier motion_modifier;
       struct gfc_common_head *common;
       bool lastprivate_conditional;
     } u;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index abca146d78e..77598190600 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -394,7 +394,8 @@  gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
 			     gfc_omp_namelist ***headp = NULL,
 			     bool allow_sections = false,
 			     bool allow_derived = false,
-			     bool *has_all_memory = NULL)
+			     bool *has_all_memory = NULL,
+			     bool allow_motion_modifier = false)
 {
   gfc_omp_namelist *head, *tail, *p;
   locus old_loc, cur_loc;
@@ -402,6 +403,7 @@  gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
   gfc_symbol *sym;
   match m;
   gfc_symtree *st;
+  bool present = false;
 
   head = tail = NULL;
 
@@ -437,6 +439,12 @@  gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
 	  tail->where = cur_loc;
 	  goto next_item;
 	}
+      else if (allow_motion_modifier && m == MATCH_YES && strcmp (n, "present") == 0
+	       && gfc_match_char (':') == MATCH_YES)
+	{
+	  present = true;
+	  m = gfc_match_name (n);
+	}
       if (m == MATCH_YES)
 	{
 	  gfc_symtree *st;
@@ -537,6 +545,13 @@  gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
   *list = head;
   if (headp)
     *headp = list;
+
+  if (present)
+    {
+      gfc_omp_namelist *n;
+      for (n = head; n; n = n->next)
+	n->u.motion_modifier = OMP_MOTION_PRESENT;
+    }
   return MATCH_YES;
 
 syntax:
@@ -2087,6 +2102,8 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		behavior = OMP_DEFAULTMAP_FROM;
 	      else if (gfc_match ("firstprivate ") == MATCH_YES)
 		behavior = OMP_DEFAULTMAP_FIRSTPRIVATE;
+	      else if (gfc_match ("present ") == MATCH_YES)
+		behavior = OMP_DEFAULTMAP_PRESENT;
 	      else if (gfc_match ("none ") == MATCH_YES)
 		behavior = OMP_DEFAULTMAP_NONE;
 	      else if (gfc_match ("default ") == MATCH_YES)
@@ -2094,7 +2111,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      else
 		{
 		  gfc_error ("Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, "
-			   "NONE or DEFAULT at %C");
+			     "PRESENT, NONE or DEFAULT at %C");
 		  break;
 		}
 	      if (')' == gfc_peek_ascii_char ())
@@ -2520,7 +2537,8 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_FROM)
 	      && (gfc_match_omp_variable_list ("from (",
 					      &c->lists[OMP_LIST_FROM], false,
-					      NULL, &head, true, true)
+					      NULL, &head, true, true, NULL,
+					      true)
 		  == MATCH_YES))
 	    continue;
 	  break;
@@ -2877,6 +2895,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      locus old_loc2 = gfc_current_locus;
 	      int always_modifier = 0;
 	      int close_modifier = 0;
+	      int present_modifier = 0;
 	      locus second_always_locus = old_loc2;
 	      locus second_close_locus = old_loc2;
 
@@ -2893,20 +2912,38 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		      if (close_modifier++ == 1)
 			second_close_locus = current_locus;
 		    }
+		  else if (gfc_match ("present ") == MATCH_YES)
+		    {
+		      if (present_modifier++ == 1)
+			second_close_locus = current_locus;
+		    }
 		  else
 		    break;
 		  gfc_match (", ");
 		}
 
 	      gfc_omp_map_op map_op = OMP_MAP_TOFROM;
+	      int always_present_modifier
+		= always_modifier && present_modifier;
+
 	      if (gfc_match ("alloc : ") == MATCH_YES)
-		map_op = OMP_MAP_ALLOC;
+		map_op = present_modifier ? OMP_MAP_PRESENT_ALLOC
+			 : OMP_MAP_ALLOC;
 	      else if (gfc_match ("tofrom : ") == MATCH_YES)
-		map_op = always_modifier ? OMP_MAP_ALWAYS_TOFROM : OMP_MAP_TOFROM;
+		map_op = always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_TOFROM
+			 : present_modifier ? OMP_MAP_PRESENT_TOFROM
+			 : always_modifier ? OMP_MAP_ALWAYS_TOFROM
+			 : OMP_MAP_TOFROM;
 	      else if (gfc_match ("to : ") == MATCH_YES)
-		map_op = always_modifier ? OMP_MAP_ALWAYS_TO : OMP_MAP_TO;
+		map_op = always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_TO
+			 : present_modifier ? OMP_MAP_PRESENT_TO
+			 : always_modifier ? OMP_MAP_ALWAYS_TO
+			 : OMP_MAP_TO;
 	      else if (gfc_match ("from : ") == MATCH_YES)
-		map_op = always_modifier ? OMP_MAP_ALWAYS_FROM : OMP_MAP_FROM;
+		map_op = always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_FROM
+			 : present_modifier ? OMP_MAP_PRESENT_FROM
+			 : always_modifier ? OMP_MAP_ALWAYS_FROM
+			 : OMP_MAP_FROM;
 	      else if (gfc_match ("release : ") == MATCH_YES)
 		map_op = OMP_MAP_RELEASE;
 	      else if (gfc_match ("delete : ") == MATCH_YES)
@@ -3458,7 +3495,8 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  else if ((mask & OMP_CLAUSE_TO)
 	      && (gfc_match_omp_variable_list ("to (",
 					      &c->lists[OMP_LIST_TO], false,
-					      NULL, &head, true, true)
+					      NULL, &head, true, true, NULL,
+					      true)
 		  == MATCH_YES))
 	    continue;
 	  break;
@@ -7805,11 +7843,18 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			{
 			case OMP_MAP_TO:
 			case OMP_MAP_ALWAYS_TO:
+			case OMP_MAP_PRESENT_TO:
+			case OMP_MAP_ALWAYS_PRESENT_TO:
 			case OMP_MAP_FROM:
 			case OMP_MAP_ALWAYS_FROM:
+			case OMP_MAP_PRESENT_FROM:
+			case OMP_MAP_ALWAYS_PRESENT_FROM:
 			case OMP_MAP_TOFROM:
 			case OMP_MAP_ALWAYS_TOFROM:
+			case OMP_MAP_PRESENT_TOFROM:
+			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
 			case OMP_MAP_ALLOC:
+			case OMP_MAP_PRESENT_ALLOC:
 			  break;
 			default:
 			  gfc_error ("TARGET%s with map-type other than TO, "
@@ -7825,6 +7870,8 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			{
 			case OMP_MAP_TO:
 			case OMP_MAP_ALWAYS_TO:
+			case OMP_MAP_PRESENT_TO:
+			case OMP_MAP_ALWAYS_PRESENT_TO:
 			case OMP_MAP_ALLOC:
 			  break;
 			case OMP_MAP_TOFROM:
@@ -7833,6 +7880,12 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			case OMP_MAP_ALWAYS_TOFROM:
 			  n->u.map_op = OMP_MAP_ALWAYS_TO;
 			  break;
+			case OMP_MAP_PRESENT_TOFROM:
+			  n->u.map_op = OMP_MAP_PRESENT_TO;
+			  break;
+			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
+			  break;
 			default:
 			  gfc_error ("TARGET ENTER DATA with map-type other "
 				     "than TO, TOFROM or ALLOC on MAP clause "
@@ -7845,6 +7898,8 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			{
 			case OMP_MAP_FROM:
 			case OMP_MAP_ALWAYS_FROM:
+			case OMP_MAP_PRESENT_FROM:
+			case OMP_MAP_ALWAYS_PRESENT_FROM:
 			case OMP_MAP_RELEASE:
 			case OMP_MAP_DELETE:
 			  break;
@@ -7854,6 +7909,12 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			case OMP_MAP_ALWAYS_TOFROM:
 			  n->u.map_op = OMP_MAP_ALWAYS_FROM;
 			  break;
+			case OMP_MAP_PRESENT_TOFROM:
+			  n->u.map_op = OMP_MAP_PRESENT_FROM;
+			  break;
+			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
+			  break;
 			default:
 			  gfc_error ("TARGET EXIT DATA with map-type other "
 				     "than FROM, TOFROM, RELEASE, or DELETE on "
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 2d16f3be8ea..35f9669c6c0 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3066,6 +3066,30 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  always_modifier = true;
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
 		  break;
+		case OMP_MAP_PRESENT_ALLOC:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_ALLOC);
+		  break;
+		case OMP_MAP_PRESENT_TO:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_TO);
+		  break;
+		case OMP_MAP_PRESENT_FROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_FROM);
+		  break;
+		case OMP_MAP_PRESENT_TOFROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_TOFROM);
+		  break;
+		case OMP_MAP_ALWAYS_PRESENT_TO:
+		  always_modifier = true;
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_TO);
+		  break;
+		case OMP_MAP_ALWAYS_PRESENT_FROM:
+		  always_modifier = true;
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_FROM);
+		  break;
+		case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+		  always_modifier = true;
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_TOFROM);
+		  break;
 		case OMP_MAP_RELEASE:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
 		  break;
@@ -3705,6 +3729,9 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
 		  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
 		}
+	      if (n->u.motion_modifier == OMP_MOTION_PRESENT)
+		OMP_CLAUSE_SET_MOTION_MODIFIER (node,
+						OMP_CLAUSE_MOTION_PRESENT);
 	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
 	    }
 	  break;
@@ -4246,6 +4273,9 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	case OMP_DEFAULTMAP_FIRSTPRIVATE:
 	  behavior = OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE;
 	  break;
+	case OMP_DEFAULTMAP_PRESENT:
+	  behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+	  break;
 	case OMP_DEFAULTMAP_NONE: behavior = OMP_CLAUSE_DEFAULTMAP_NONE; break;
 	case OMP_DEFAULTMAP_DEFAULT:
 	  behavior = OMP_CLAUSE_DEFAULTMAP_DEFAULT;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 96845154a92..10495fbd77e 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7904,6 +7904,11 @@  omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 		  else if (ctx->defaultmap[gdmk]
 			   & (GOVD_MAP_0LEN_ARRAY | GOVD_FIRSTPRIVATE))
 		    nflags |= ctx->defaultmap[gdmk];
+		  else if (ctx->defaultmap[gdmk] & GOVD_MAP_FORCE_PRESENT)
+		    {
+		      gcc_assert (ctx->defaultmap[gdmk] & GOVD_MAP);
+		      nflags |= ctx->defaultmap[gdmk] | GOVD_MAP_ALLOC_ONLY;
+		    }
 		  else
 		    {
 		      gcc_assert (ctx->defaultmap[gdmk] & GOVD_MAP);
@@ -9062,6 +9067,13 @@  omp_get_attachment (omp_mapping_group *grp)
     case GOMP_MAP_FORCE_TO:
     case GOMP_MAP_FORCE_TOFROM:
     case GOMP_MAP_FORCE_PRESENT:
+    case GOMP_MAP_PRESENT_ALLOC:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_PRESENT_TO:
+    case GOMP_MAP_PRESENT_TOFROM:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_TO:
+    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
     case GOMP_MAP_ALLOC:
     case GOMP_MAP_RELEASE:
     case GOMP_MAP_DELETE:
@@ -9293,6 +9305,13 @@  omp_group_base (omp_mapping_group *grp, unsigned int *chained,
     case GOMP_MAP_FORCE_TO:
     case GOMP_MAP_FORCE_TOFROM:
     case GOMP_MAP_FORCE_PRESENT:
+    case GOMP_MAP_PRESENT_ALLOC:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_PRESENT_TO:
+    case GOMP_MAP_PRESENT_TOFROM:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_TO:
+    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
     case GOMP_MAP_ALLOC:
     case GOMP_MAP_RELEASE:
     case GOMP_MAP_DELETE:
@@ -10776,6 +10795,50 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  delete grpmap;
 	  delete groups;
 	}
+
+      /* OpenMP map clauses with 'present' need to go in front of those
+	 without.  */
+      tree present_map_head = NULL;
+      tree *present_map_tail_p = &present_map_head;
+      tree *first_map_clause_p = NULL;
+
+      for (tree *c_p = list_p; *c_p; )
+	{
+	  tree c = *c_p;
+	  tree *next_c_p = &OMP_CLAUSE_CHAIN (c);
+
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	    {
+	      if (!first_map_clause_p)
+		first_map_clause_p = c_p;
+	      switch (OMP_CLAUSE_MAP_KIND (c))
+		{
+		case GOMP_MAP_PRESENT_ALLOC:
+		case GOMP_MAP_PRESENT_FROM:
+		case GOMP_MAP_PRESENT_TO:
+		case GOMP_MAP_PRESENT_TOFROM:
+		  next_c_p = c_p;
+		  *c_p = OMP_CLAUSE_CHAIN (c);
+
+		  OMP_CLAUSE_CHAIN (c) = NULL;
+		  *present_map_tail_p = c;
+		  present_map_tail_p = &OMP_CLAUSE_CHAIN (c);
+
+		  break;
+
+		default:
+		  break;
+		}
+	    }
+
+	  c_p = next_c_p;
+	}
+      if (first_map_clause_p && present_map_head)
+	{
+	  tree next = *first_map_clause_p;
+	  *first_map_clause_p = present_map_head;
+	  *present_map_tail_p = next;
+	}
     }
   else if (region_type & ORT_ACC)
     {
@@ -11972,6 +12035,9 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      case OMP_CLAUSE_DEFAULTMAP_NONE:
 		ctx->defaultmap[gdmk] = 0;
 		break;
+	      case OMP_CLAUSE_DEFAULTMAP_PRESENT:
+		ctx->defaultmap[gdmk] = GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+		break;
 	      case OMP_CLAUSE_DEFAULTMAP_DEFAULT:
 		switch (gdmk)
 		  {
@@ -12416,6 +12482,9 @@  gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 	case GOVD_MAP_FORCE_PRESENT:
 	  kind = GOMP_MAP_FORCE_PRESENT;
 	  break;
+	case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY:
+	  kind = GOMP_MAP_PRESENT_ALLOC;
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index fef41a013ec..c1ac34e5ec0 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1576,6 +1576,9 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_TO
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_FROM
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_TOFROM
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable
@@ -12797,6 +12800,14 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_ALLOC:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_STRUCT:
@@ -13338,6 +13349,13 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    case GOMP_MAP_ALWAYS_TO:
 		    case GOMP_MAP_ALWAYS_FROM:
 		    case GOMP_MAP_ALWAYS_TOFROM:
+		    case GOMP_MAP_PRESENT_ALLOC:
+		    case GOMP_MAP_PRESENT_TO:
+		    case GOMP_MAP_PRESENT_FROM:
+		    case GOMP_MAP_PRESENT_TOFROM:
+		    case GOMP_MAP_ALWAYS_PRESENT_TO:
+		    case GOMP_MAP_ALWAYS_PRESENT_FROM:
+		    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 		    case GOMP_MAP_RELEASE:
 		    case GOMP_MAP_FORCE_TO:
 		    case GOMP_MAP_FORCE_FROM:
@@ -13377,11 +13395,15 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		tkind_zero = tkind;
 		break;
 	      case OMP_CLAUSE_TO:
-		tkind = GOMP_MAP_TO;
+		tkind
+		  = OMP_CLAUSE_MOTION_MODIFIER (c) == OMP_CLAUSE_MOTION_PRESENT
+		    ? GOMP_MAP_PRESENT_TO : GOMP_MAP_TO;
 		tkind_zero = tkind;
 		break;
 	      case OMP_CLAUSE_FROM:
-		tkind = GOMP_MAP_FROM;
+		tkind
+		  = OMP_CLAUSE_MOTION_MODIFIER (c) == OMP_CLAUSE_MOTION_PRESENT
+		    ? GOMP_MAP_PRESENT_FROM : GOMP_MAP_FROM;
 		tkind_zero = tkind;
 		break;
 	      default:
diff --git a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
new file mode 100644
index 00000000000..1afff7ea38f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
@@ -0,0 +1,24 @@ 
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+  int a[N], b[N], c[N];
+
+  /* Should generate implicit 'map(present, alloc)' clauses.  */
+  #pragma omp target defaultmap (present: aggregate)
+    for (int i = 0; i < N; i++)
+      c[i] = a[i] + b[i];
+
+  /* Should generate implicit 'map(present, alloc)' clauses,
+     and they should go before other non-present clauses.  */
+  #pragma omp target map(from: c) defaultmap (present: aggregate)
+    for (int i = 0; i < N; i++)
+      c[i] = a[i] + b[i];
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(present,alloc:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 6ee59714847..b4683ddbabf 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,10 +13,10 @@  foo (void)
   #pragma omp target map (to:a)
   ;
 
-  #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+  #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'present'" } */
   ;
 
-  #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+  #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'present'" } */
   ;
 
   #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
diff --git a/gcc/testsuite/c-c++-common/gomp/map-8.c b/gcc/testsuite/c-c++-common/gomp/map-8.c
new file mode 100644
index 00000000000..4b4bd6d2aa3
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-8.c
@@ -0,0 +1,32 @@ 
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+  int a[N], b[N], c[N];
+
+  /* Should be able to parse 'present' map modifier.  */
+  #pragma omp target enter data map (present, to: a, b)
+
+  #pragma omp target data map (present, to: a, b) map (always, present, from: c)
+
+  #pragma omp target map (present, to: a, b) map (present, from: c)
+    for (int i = 0; i < N; i++)
+      c[i] = a[i] + b[i];
+
+  #pragma omp target exit data map (always, present, from: c)
+
+  /* Map clauses with 'present' modifier should go ahead of those without.  */
+  #pragma omp target map (to: a) map (present, to: b) map (from: c)
+    for (int i = 0; i < N; i++)
+      c[i] = a[i] + b[i];
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target data map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*map\\(present,from:c \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-1.c
new file mode 100644
index 00000000000..0233fe5a7af
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-1.c
@@ -0,0 +1,15 @@ 
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+  int a[N], b[N];
+
+  /* Should be able to parse present in to/from clauses of 'target update'.  */
+  #pragma omp target update to(present: a) from(present: b)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target update from\\(present:b \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
index 299d971f23c..1f1b8528aef 100644
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
@@ -2,7 +2,7 @@ 
 
 implicit none
 
-!$omp target defaultmap(bar)  ! { dg-error "25: Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, NONE or DEFAULT" }
+!$omp target defaultmap(bar)  ! { dg-error "25: Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, PRESENT, NONE or DEFAULT" }
 
 !$omp target defaultmap ( alloc: foo)  ! { dg-error "34: Expected SCALAR, AGGREGATE, ALLOCATABLE or POINTER" }
 
diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
new file mode 100644
index 00000000000..669a623f746
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
@@ -0,0 +1,26 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+  implicit none
+  integer, parameter :: N = 1000
+  integer :: a(N), b(N), c(N), i
+  
+  ! Should generate implicit 'map(present, alloc)' clauses.
+  !$omp target defaultmap (present: aggregate)
+    do i = 1, N
+      c(i) = a(i) + b(i)
+    end do
+  !$omp end target
+
+  ! Should generate implicit 'map(present, alloc)' clauses,
+  ! and they should go before other non-present clauses.
+  !$omp target map(from: c) defaultmap (present: aggregate)
+    do i = 1, N
+      c(i) = a(i) + b(i)
+    end do
+  !$omp end target
+end program
+  
+! { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\).*map\\(present,alloc:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) defaultmap\\(present:aggregate\\)" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 b/gcc/testsuite/gfortran.dg/gomp/map-9.f90
new file mode 100644
index 00000000000..cc87212f8d0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-9.f90
@@ -0,0 +1,34 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+  implicit none
+  integer, parameter :: N = 1000
+  integer :: a(N), b(N), c(N), i
+
+  ! Should be able to parse 'present' map modifier.
+  !$omp target enter data map (present, to: a, b)
+
+  !$omp target data map (present, to: a, b) map (always, present, from: c)
+    !$omp target map (present, to: a, b) map (present, from: c)
+      do i = 1, N
+	c(i) = a(i) + b(i)
+      end do
+    !$omp end target
+  !$omp end target data
+
+  !$omp target exit data map (always, present, from: c)
+
+  ! Map clauses with 'present' modifier should go ahead of those without.
+  !$omp target map (to: a) map (present, to: b) map (from: c)
+    do i = 1, N
+      c(i) = a(i) + b(i)
+    end do
+  !$omp end target
+end program
+
+! { dg-final { scan-tree-dump "pragma omp target enter data map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target data map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target exit data map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
new file mode 100644
index 00000000000..a382b87f229
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
@@ -0,0 +1,13 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+  implicit none
+  integer, parameter :: N = 1000
+  integer :: a(N), b(N), i
+
+  ! Should be able to parse present in to/from clauses of 'target update'.
+  !$omp target update to(present: a) from(present: b)
+end program
+
+! { dg-final { scan-tree-dump "pragma omp target update to\\(present:a \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index acd8deea34e..e50df80b22c 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -573,7 +573,8 @@  enum omp_clause_defaultmap_kind {
   OMP_CLAUSE_DEFAULTMAP_NONE = 6 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
   OMP_CLAUSE_DEFAULTMAP_DEFAULT
     = 7 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
-  OMP_CLAUSE_DEFAULTMAP_MASK = 7 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1)
+  OMP_CLAUSE_DEFAULTMAP_PRESENT = 8 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
+  OMP_CLAUSE_DEFAULTMAP_MASK = 15 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1)
 };
 
 enum omp_clause_bind_kind {
@@ -582,6 +583,11 @@  enum omp_clause_bind_kind {
   OMP_CLAUSE_BIND_THREAD
 };
 
+enum omp_clause_motion_modifier {
+  OMP_CLAUSE_MOTION_NONE,
+  OMP_CLAUSE_MOTION_PRESENT
+};
+
 /* memory-order-clause on OpenMP atomic/flush constructs or
    argument of atomic_default_mem_order clause.  */
 enum omp_memory_order {
@@ -1646,6 +1652,7 @@  struct GTY(()) tree_omp_clause {
     enum omp_clause_defaultmap_kind defaultmap_kind;
     enum omp_clause_bind_kind      bind_kind;
     enum omp_clause_device_type_kind device_type_kind;
+    enum omp_clause_motion_modifier motion_modifier;
   } GTY ((skip)) subcode;
 
   /* The gimplification of OMP_CLAUSE_REDUCTION_{INIT,MERGE} for omp-low's
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 7947f9647a1..99f8acc3acc 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -991,6 +991,27 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	  pp_string (pp, "attach_zero_length_array_section");
 	  break;
+	case GOMP_MAP_PRESENT_ALLOC:
+	  pp_string (pp, "present,alloc");
+	  break;
+	case GOMP_MAP_PRESENT_TO:
+	  pp_string (pp, "present,to");
+	  break;
+	case GOMP_MAP_PRESENT_FROM:
+	  pp_string (pp, "present,from");
+	  break;
+	case GOMP_MAP_PRESENT_TOFROM:
+	  pp_string (pp, "present,tofrom");
+	  break;
+	case GOMP_MAP_ALWAYS_PRESENT_TO:
+	  pp_string (pp, "always,present,to");
+	  break;
+	case GOMP_MAP_ALWAYS_PRESENT_FROM:
+	  pp_string (pp, "always,present,from");
+	  break;
+	case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	  pp_string (pp, "always,present,tofrom");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -1038,12 +1059,16 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     case OMP_CLAUSE_FROM:
       pp_string (pp, "from(");
+      if (OMP_CLAUSE_MOTION_MODIFIER (clause) == OMP_CLAUSE_MOTION_PRESENT)
+	pp_string (pp, "present:");
       dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
 			 spc, flags, false);
       goto print_clause_size;
 
     case OMP_CLAUSE_TO:
       pp_string (pp, "to(");
+      if (OMP_CLAUSE_MOTION_MODIFIER (clause) == OMP_CLAUSE_MOTION_PRESENT)
+	pp_string (pp, "present:");
       dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
 			 spc, flags, false);
       goto print_clause_size;
@@ -1210,6 +1235,9 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case OMP_CLAUSE_DEFAULTMAP_NONE:
 	  pp_string (pp, "none");
 	  break;
+	case OMP_CLAUSE_DEFAULTMAP_PRESENT:
+	  pp_string (pp, "present");
+	  break;
 	case OMP_CLAUSE_DEFAULTMAP_DEFAULT:
 	  pp_string (pp, "default");
 	  break;
diff --git a/gcc/tree.h b/gcc/tree.h
index c656cd5b7bf..eb573a01b55 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1742,6 +1742,12 @@  class auto_suppress_location_wrappers
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \
    = (unsigned int) (MAP_KIND))
 
+#define OMP_CLAUSE_MOTION_MODIFIER(NODE) \
+  ((enum omp_clause_motion_modifier) OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->omp_clause.subcode.motion_modifier)
+#define OMP_CLAUSE_SET_MOTION_MODIFIER(NODE, MOTION_MODIFIER) \
+  (OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->omp_clause.subcode.motion_modifier \
+   = (MOTION_MODIFIER))
+
 /* Nonzero if this map clause is for array (rather than pointer) based array
    section with zero bias.  Both the non-decl OMP_CLAUSE_MAP and corresponding
    OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag.  */
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 1b9b07dc245..49b7dd86ff5 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -42,6 +42,7 @@ 
 #define GOMP_MAP_FLAG_SPECIAL_2		(1 << 4)
 #define GOMP_MAP_FLAG_SPECIAL_3		(1 << 5)
 #define GOMP_MAP_FLAG_SPECIAL_4		(1 << 6)
+#define GOMP_MAP_FLAG_SPECIAL_5		(1 << 7)
 #define GOMP_MAP_FLAG_SPECIAL		(GOMP_MAP_FLAG_SPECIAL_1 \
 					 | GOMP_MAP_FLAG_SPECIAL_0)
 #define GOMP_MAP_DEEP_COPY		(GOMP_MAP_FLAG_SPECIAL_4 \
@@ -55,9 +56,14 @@ 
 					 | GOMP_MAP_FLAG_SPECIAL_1 \
 					 | GOMP_MAP_FLAG_SPECIAL_2 \
 					 | GOMP_MAP_FLAG_SPECIAL_3 \
-					 | GOMP_MAP_FLAG_SPECIAL_4)
+					 | GOMP_MAP_FLAG_SPECIAL_4 \
+					 | GOMP_MAP_FLAG_SPECIAL_5)
 /* Flag to force a specific behavior (or else, trigger a run-time error).  */
-#define GOMP_MAP_FLAG_FORCE		(1 << 7)
+#define GOMP_MAP_FLAG_FORCE		(GOMP_MAP_FLAG_SPECIAL_5)
+#define GOMP_MAP_FLAG_PRESENT		(GOMP_MAP_FLAG_SPECIAL_5 \
+					 | GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_FLAG_ALWAYS_PRESENT	(GOMP_MAP_FLAG_SPECIAL_2 \
+					 | GOMP_MAP_FLAG_PRESENT)
 
 enum gomp_map_kind
   {
@@ -130,6 +136,23 @@  enum gomp_map_kind
        device.  */
     GOMP_MAP_ALWAYS_TOFROM =		(GOMP_MAP_FLAG_SPECIAL_2
 					 | GOMP_MAP_TOFROM),
+    /* Must already be present.  */
+    GOMP_MAP_PRESENT_ALLOC =		(GOMP_MAP_FLAG_PRESENT | GOMP_MAP_ALLOC),
+    /* Must already be present, copy to device.  */
+    GOMP_MAP_PRESENT_TO =		(GOMP_MAP_FLAG_PRESENT | GOMP_MAP_TO),
+    /* Must already be present, copy from device.  */
+    GOMP_MAP_PRESENT_FROM =		(GOMP_MAP_FLAG_PRESENT | GOMP_MAP_FROM),
+    /* Must already be present, copy to and from device.  */
+    GOMP_MAP_PRESENT_TOFROM =		(GOMP_MAP_FLAG_PRESENT | GOMP_MAP_TOFROM),
+    /* Must already be present, unconditionally copy to device.  */
+    GOMP_MAP_ALWAYS_PRESENT_TO =	(GOMP_MAP_FLAG_ALWAYS_PRESENT
+					 | GOMP_MAP_TO),
+    /* Must already be present, unconditionally copy from device.  */
+    GOMP_MAP_ALWAYS_PRESENT_FROM =	(GOMP_MAP_FLAG_ALWAYS_PRESENT
+					 | GOMP_MAP_FROM),
+    /* Must already be present, unconditionally copy to and from device.  */
+    GOMP_MAP_ALWAYS_PRESENT_TOFROM =	(GOMP_MAP_FLAG_ALWAYS_PRESENT
+					 | GOMP_MAP_TOFROM),
     /* Map a sparse struct; the address is the base of the structure, alignment
        it's required alignment, and size is the number of adjacent entries
        that belong to the struct.  The adjacent entries should be sorted by
@@ -186,11 +209,11 @@  enum gomp_map_kind
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
-  (!((X) & GOMP_MAP_FLAG_SPECIAL) \
+  ((!((X) & GOMP_MAP_FLAG_SPECIAL) || GOMP_MAP_PRESENT_P (X)) \
    && ((X) & GOMP_MAP_FLAG_TO))
 
 #define GOMP_MAP_COPY_FROM_P(X) \
-  (!((X) & GOMP_MAP_FLAG_SPECIAL) \
+  ((!((X) & GOMP_MAP_FLAG_SPECIAL) || GOMP_MAP_PRESENT_P (X)) \
    && ((X) & GOMP_MAP_FLAG_FROM))
 
 #define GOMP_MAP_ALWAYS_POINTER_P(X) \
@@ -201,17 +224,27 @@  enum gomp_map_kind
    || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
 
 #define GOMP_MAP_ALWAYS_TO_P(X) \
-  (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+  (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM) \
+   || ((X) == GOMP_MAP_ALWAYS_PRESENT_TO) \
+   || ((X) == GOMP_MAP_ALWAYS_PRESENT_TOFROM))
 
 #define GOMP_MAP_ALWAYS_FROM_P(X) \
-  (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+  (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM) \
+   || ((X) == GOMP_MAP_ALWAYS_PRESENT_FROM) \
+   || ((X) == GOMP_MAP_ALWAYS_PRESENT_TOFROM))
 
 #define GOMP_MAP_ALWAYS_P(X) \
-  (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
+  (GOMP_MAP_ALWAYS_TO_P (X) || GOMP_MAP_ALWAYS_FROM_P (X))
 
 #define GOMP_MAP_IMPLICIT_P(X) \
   (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT)
 
+#define GOMP_MAP_FORCE_P(X) \
+  (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_FLAG_FORCE)
+
+#define GOMP_MAP_PRESENT_P(X) \
+  (((X) & GOMP_MAP_FLAG_PRESENT) == GOMP_MAP_FLAG_PRESENT)
+
 
 /* Asynchronous behavior.  Keep in sync with
    libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t.  */
diff --git a/libgomp/target.c b/libgomp/target.c
index 483851c95ac..358a1c267f7 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -360,6 +360,8 @@  gomp_to_device_kind_p (int kind)
     case GOMP_MAP_FORCE_ALLOC:
     case GOMP_MAP_FORCE_FROM:
     case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
       return false;
     default:
       return true;
@@ -592,7 +594,7 @@  gomp_map_vars_existing (struct gomp_device_descr *devicep,
   else
     tgt_var->length = newn->host_end - newn->host_start;
 
-  if ((kind & GOMP_MAP_FLAG_FORCE)
+  if (GOMP_MAP_FORCE_P (kind)
       /* For implicit maps, old contained in new is valid.  */
       || !(implicit_subset
 	   /* Otherwise, new contained inside old is considered valid.  */
@@ -1708,6 +1710,20 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 #endif
 		    }
 		    break;
+		  case GOMP_MAP_PRESENT_ALLOC:
+		  case GOMP_MAP_PRESENT_TO:
+		  case GOMP_MAP_PRESENT_FROM:
+		  case GOMP_MAP_PRESENT_TOFROM:
+		  case GOMP_MAP_ALWAYS_PRESENT_TO:
+		  case GOMP_MAP_ALWAYS_PRESENT_FROM:
+		  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+		    /* We already looked up the memory region above and it
+		       was missing.  */
+		    gomp_mutex_unlock (&devicep->lock);
+		    gomp_fatal ("present clause: !omp_target_is_present "
+				"(%p, %d)",
+				(void *) k->host_start, devicep->target_id);
+		    break;
 		  case GOMP_MAP_FORCE_DEVICEPTR:
 		    assert (k->host_end - k->host_start == sizeof (void *));
 		    gomp_copy_host2dev (devicep, aq,
@@ -2117,6 +2133,20 @@  gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 		  gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
 	      }
 	  }
+	else
+	  {
+	    int kind = get_kind (short_mapkind, kinds, i);
+
+	    if (GOMP_MAP_PRESENT_P (kind))
+	      {
+		/* We already looked up the memory region above and it
+		   was missing.  */
+		gomp_mutex_unlock (&devicep->lock);
+		gomp_fatal ("present clause: !omp_target_is_present "
+			    "(%p, %d)",
+			    (void *) hostaddrs[i], devicep->target_id);
+	      }
+	  }
       }
   gomp_mutex_unlock (&devicep->lock);
 }
@@ -3425,7 +3455,8 @@  gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_DELETE:
 	      case GOMP_MAP_RELEASE:
 	      case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
-		/* Assume it is present; look it up - but ignore otherwise. */
+		/* Assume it is present; look it up - but ignore unless the
+		   present clause is there. */
 	      case GOMP_MAP_ALLOC:
 	      case GOMP_MAP_FROM:
 	      case GOMP_MAP_FORCE_ALLOC:
@@ -3437,6 +3468,12 @@  gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_FORCE_TOFROM:
 	      case GOMP_MAP_ALWAYS_TO:
 	      case GOMP_MAP_ALWAYS_TOFROM:
+	      case GOMP_MAP_PRESENT_FROM:
+	      case GOMP_MAP_PRESENT_TO:
+	      case GOMP_MAP_PRESENT_TOFROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_FROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_TO:
+	      case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
 		cdata[i].devaddr = devaddrs[i];
 		bool zero_len = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
@@ -3457,7 +3494,23 @@  gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 					      devaddrs[i] + sizes[i], zero_len);
 		    cdata[i].present = n2 != NULL;
 		  }
-		if (!cdata[i].present
+		if (!cdata[i].present && GOMP_MAP_PRESENT_P (kind))
+		  {
+		    gomp_mutex_unlock (&devicep->lock);
+#ifdef HAVE_INTTYPES_H
+		    gomp_fatal ("present clause: no corresponding data on "
+				"parent device at %p with size %"PRIu64,
+				(void *) (uintptr_t) devaddrs[i],
+				(uint64_t) sizes[i]);
+#else
+		    gomp_fatal ("present clause: no corresponding data on "
+				"parent device at %p with size %lu",
+				(void *) (uintptr_t) devaddrs[i],
+				(unsigned long) sizes[i]);
+#endif
+		    break;
+		  }
+		else if (!cdata[i].present
 		    && kind != GOMP_MAP_DELETE
 		    && kind != GOMP_MAP_RELEASE
 		    && kind != GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION)
@@ -3475,8 +3528,7 @@  gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 		     && (kind == GOMP_MAP_TO || kind == GOMP_MAP_TOFROM))
 		    || kind == GOMP_MAP_FORCE_TO
 		    || kind == GOMP_MAP_FORCE_TOFROM
-		    || kind == GOMP_MAP_ALWAYS_TO
-		    || kind == GOMP_MAP_ALWAYS_TOFROM)
+		    || GOMP_MAP_ALWAYS_TO_P (kind))
 		  {
 		    if (dev_to_host_cpy)
 		      dev_to_host_cpy ((void *) (uintptr_t) devaddrs[i],
@@ -3661,6 +3713,10 @@  gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_FORCE_TOFROM:
 	      case GOMP_MAP_ALWAYS_FROM:
 	      case GOMP_MAP_ALWAYS_TOFROM:
+	      case GOMP_MAP_PRESENT_FROM:
+	      case GOMP_MAP_PRESENT_TOFROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_FROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 		copy = true;
 		/* FALLTHRU */
 	      case GOMP_MAP_FROM:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
new file mode 100644
index 00000000000..bbc4559b12e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
@@ -0,0 +1,27 @@ 
+/* { dg-do run { target offload_target_any } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+  int a[N], b[N], c[N];
+
+  for (int i = 0; i < N; i++) {
+    a[i] = i * 2;
+    b[i] = i * 3 + 1;
+  }
+
+  #pragma omp target enter data map (alloc: a, c)
+    /* a has already been allocated, so this should be okay.  */
+    #pragma omp target map (present, to: a)
+      for (int i = 0; i < N; i++)
+	c[i] = a[i];
+
+    /* b has not been allocated, so this should result in an error.  */
+    /* { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+    #pragma omp target map (present, to: b)
+      for (int i = 0; i < N; i++)
+	c[i] += b[i];
+  #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
new file mode 100644
index 00000000000..6259c959c04
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
@@ -0,0 +1,27 @@ 
+/* { dg-do run { target offload_target_any } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+  int a[N], b[N], c[N];
+
+  for (int i = 0; i < N; i++) {
+    a[i] = i * 2;
+    b[i] = i * 3 + 1;
+  }
+
+  #pragma omp target enter data map (alloc: a, c)
+    /* a has already been allocated, so this should be okay.  */
+    #pragma omp target defaultmap (present)
+      for (int i = 0; i < N; i++)
+	c[i] = a[i];
+
+    /* b has not been allocated, so this should result in an error.  */
+    /* { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+    #pragma omp target defaultmap (present)
+      for (int i = 0; i < N; i++)
+	c[i] += b[i];
+  #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
new file mode 100644
index 00000000000..89e648645b2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
@@ -0,0 +1,27 @@ 
+/* { dg-do run { target offload_target_any } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#include <stdio.h>
+
+#define N 100
+
+int main (void)
+{
+  int a[N], b[N], c[N];
+
+  for (int i = 0; i < N; i++) {
+    a[i] = i * 2;
+    b[i] = i * 3 + 1;
+  }
+
+  #pragma omp target enter data map (alloc: a, c)
+
+  /* This should work as a has already been allocated.  */
+  #pragma omp target update to (present: a)
+
+  /* This should fail as b has not been allocated.  */
+  /* { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+  #pragma omp target update to (present: b)
+
+  #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
new file mode 100644
index 00000000000..80046011b25
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
@@ -0,0 +1,30 @@ 
+! { dg-do run { target offload_target_any } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N), b(N), c(N), i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3 + 1
+  end do
+
+  !$omp target enter data map (alloc: a)
+    ! a has already been allocated, so this should be okay.
+    !$omp target map (present, to: a)
+      do i = 1, N
+	c(i) = a(i)
+      end do
+    !$omp end target
+
+    ! b has not been allocated, so this should result in an error.
+    ! { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+    !$omp target map (present, to: b)
+      do i = 1, N
+	c(i) = c(i) + b(i)
+      end do
+    !$omp end target
+  !$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
new file mode 100644
index 00000000000..0a38dea1e41
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
@@ -0,0 +1,30 @@ 
+! { dg-do run { target offload_target_any } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N), b(N), c(N), i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3 + 1
+  end do
+
+  !$omp target enter data map (alloc: a)
+    ! a has already been allocated, so this should be okay.
+    !$omp target defaultmap (present)
+      do i = 1, N
+	c(i) = a(i)
+      end do
+    !$omp end target
+
+    ! b has not been allocated, so this should result in an error.
+    ! { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+    !$omp target defaultmap (present)
+      do i = 1, N
+	c(i) = c(i) + b(i)
+      end do
+    !$omp end target
+!$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
new file mode 100644
index 00000000000..c4deb8652d1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
@@ -0,0 +1,22 @@ 
+! { dg-do run { target offload_target_any } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N), b(N), c(N), i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3 + 1
+  end do
+
+  !$omp target enter data map (alloc: a, c)
+    ! This should work as a has already been allocated.
+    !$omp target update to (present: a)
+
+    ! This should fail as b has not been allocated.
+    ! { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+    !$omp target update to (present: b)
+  !$omp target exit data map (from: c)
+end program