--- gcc/gimplify.c.jj	2015-07-14 14:29:49.274536887 +0200
+++ gcc/gimplify.c	2015-07-14 16:17:52.993914085 +0200
@@ -6394,6 +6394,13 @@ gimplify_scan_omp_clauses (tree *list_p,
 	    }
 	  goto do_notice;
 
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
+	  goto do_add;
+	case OMP_CLAUSE_IS_DEVICE_PTR:
+	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
+	  goto do_add;
+
 	do_add:
 	  decl = OMP_CLAUSE_DECL (c);
 	do_add_decl:
@@ -6957,6 +6964,8 @@ gimplify_adjust_omp_clauses (gimple_seq
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_HINT:
 	case OMP_CLAUSE_DEFAULTMAP:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
--- gcc/tree.c.jj	2015-07-14 14:29:49.032547791 +0200
+++ gcc/tree.c	2015-07-14 14:49:57.666673189 +0200
@@ -291,6 +291,8 @@ unsigned const char omp_clause_num_ops[]
   2, /* OMP_CLAUSE_FROM  */
   2, /* OMP_CLAUSE_TO  */
   2, /* OMP_CLAUSE_MAP  */
+  1, /* OMP_CLAUSE_USE_DEVICE_PTR  */
+  1, /* OMP_CLAUSE_IS_DEVICE_PTR  */
   2, /* OMP_CLAUSE__CACHE_  */
   1, /* OMP_CLAUSE_DEVICE_RESIDENT  */
   1, /* OMP_CLAUSE_USE_DEVICE  */
@@ -358,6 +360,8 @@ const char * const omp_clause_code_name[
   "from",
   "to",
   "map",
+  "use_device_ptr",
+  "is_device_ptr",
   "_cache_",
   "device_resident",
   "use_device",
@@ -11388,6 +11392,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func
 	case OMP_CLAUSE_GRAINSIZE:
 	case OMP_CLAUSE_NUM_TASKS:
 	case OMP_CLAUSE_HINT:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE__LOOPTEMP_:
 	case OMP_CLAUSE__SIMDUID_:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
--- gcc/tree-nested.c.jj	2015-07-14 14:29:49.329534408 +0200
+++ gcc/tree-nested.c	2015-07-14 14:49:57.667673145 +0200
@@ -1098,6 +1098,8 @@ convert_nonlocal_omp_clauses (tree *pcla
 	case OMP_CLAUSE_FIRSTPRIVATE:
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_SHARED:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	do_decl_clause:
 	  decl = OMP_CLAUSE_DECL (clause);
 	  if (TREE_CODE (decl) == VAR_DECL
@@ -1743,6 +1745,8 @@ convert_local_omp_clauses (tree *pclause
 	case OMP_CLAUSE_FIRSTPRIVATE:
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_SHARED:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	do_decl_clause:
 	  decl = OMP_CLAUSE_DECL (clause);
 	  if (TREE_CODE (decl) == VAR_DECL
--- gcc/tree-core.h.jj	2015-07-14 14:29:49.061546484 +0200
+++ gcc/tree-core.h	2015-07-14 14:49:57.675672794 +0200
@@ -270,6 +270,12 @@ enum omp_clause_code {
      OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list).  */
   OMP_CLAUSE_MAP,
 
+  /* OpenMP clause: use_device_ptr (variable-list).  */
+  OMP_CLAUSE_USE_DEVICE_PTR,
+
+  /* OpenMP clause: is_device_ptr (variable-list).  */
+  OMP_CLAUSE_IS_DEVICE_PTR,
+
   /* Internal structure to hold OpenACC cache directive's variable-list.
      #pragma acc cache (variable-list).  */
   OMP_CLAUSE__CACHE_,
--- gcc/omp-low.c.jj	2015-07-14 14:30:06.000000000 +0200
+++ gcc/omp-low.c	2015-07-14 16:27:54.944892284 +0200
@@ -1954,6 +1954,11 @@ scan_sharing_clauses (tree clauses, omp_
 	    }
 	  break;
 
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
+	  decl = OMP_CLAUSE_DECL (c);
+	  goto do_private;
+
 	case OMP_CLAUSE__LOOPTEMP_:
 	  gcc_assert (is_taskreg_ctx (ctx));
 	  decl = OMP_CLAUSE_DECL (c);
@@ -2138,6 +2143,8 @@ scan_sharing_clauses (tree clauses, omp_
 	  /* FALLTHRU */
 	case OMP_CLAUSE_PRIVATE:
 	case OMP_CLAUSE_LINEAR:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (is_variable_sized (decl))
 	    install_var_local (decl, ctx);
--- gcc/tree-pretty-print.c.jj	2015-07-14 14:29:49.022548242 +0200
+++ gcc/tree-pretty-print.c	2015-07-14 14:49:57.676672750 +0200
@@ -329,6 +329,12 @@ dump_omp_clause (pretty_printer *pp, tre
     case OMP_CLAUSE_UNIFORM:
       name = "uniform";
       goto print_remap;
+    case OMP_CLAUSE_USE_DEVICE_PTR:
+      name = "use_device_ptr";
+      goto print_remap;
+    case OMP_CLAUSE_IS_DEVICE_PTR:
+      name = "is_device_ptr";
+      goto print_remap;
     case OMP_CLAUSE__LOOPTEMP_:
       name = "_looptemp_";
       goto print_remap;
--- gcc/c-family/c-omp.c.jj	2015-07-14 14:29:49.281536571 +0200
+++ gcc/c-family/c-omp.c	2015-07-14 14:49:57.674672838 +0200
@@ -684,7 +684,7 @@ c_finish_omp_for (location_t locus, enum
     }
 }
 
-/* Right now we have 15 different combined constructs, this
+/* Right now we have 15 different combined/composite constructs, this
    function attempts to split or duplicate clauses for combined
    constructs.  CODE is the innermost construct in the combined construct,
    and MASK allows to determine which constructs are combined together,
@@ -744,6 +744,7 @@ c_omp_split_clauses (location_t loc, enu
 	/* First the clauses that are unique to some constructs.  */
 	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_MAP:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  s = C_OMP_CLAUSE_SPLIT_TARGET;
 	  break;
 	case OMP_CLAUSE_NUM_TEAMS:
@@ -814,7 +815,7 @@ c_omp_split_clauses (location_t loc, enu
 	  else
 	    s = C_OMP_CLAUSE_SPLIT_DISTRIBUTE;
 	  break;
-	/* Private clause is supported on all constructs but target,
+	/* Private clause is supported on all constructs,
 	   it is enough to put it on the innermost one.  For
 	   #pragma omp {for,sections} put it on parallel though,
 	   as that's what we did for OpenMP 3.1.  */
@@ -830,9 +831,18 @@ c_omp_split_clauses (location_t loc, enu
 	    }
 	  break;
 	/* Firstprivate clause is supported on all constructs but
-	   target and simd.  Put it on the outermost of those and
-	   duplicate on parallel.  */
+	   simd.  Put it on the outermost of those and duplicate on teams
+	   and parallel.  */
 	case OMP_CLAUSE_FIRSTPRIVATE:
+	  if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP))
+	      != 0)
+	    {
+	      c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
+				    OMP_CLAUSE_FIRSTPRIVATE);
+	      OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses);
+	      OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+	      cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = c;
+	    }
 	  if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS))
 	      != 0)
 	    {
--- gcc/c-family/c-pragma.h.jj	2015-07-14 14:29:49.286536346 +0200
+++ gcc/c-family/c-pragma.h	2015-07-14 14:49:57.674672838 +0200
@@ -98,6 +98,7 @@ typedef enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_HINT,
   PRAGMA_OMP_CLAUSE_IF,
   PRAGMA_OMP_CLAUSE_INBRANCH,
+  PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR,
   PRAGMA_OMP_CLAUSE_LASTPRIVATE,
   PRAGMA_OMP_CLAUSE_LINEAR,
   PRAGMA_OMP_CLAUSE_MAP,
@@ -126,6 +127,7 @@ typedef enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_TO,
   PRAGMA_OMP_CLAUSE_UNIFORM,
   PRAGMA_OMP_CLAUSE_UNTIED,
+  PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
 
   /* Clauses for Cilk Plus SIMD-enabled function.  */
   PRAGMA_CILK_CLAUSE_NOMASK,
--- gcc/c/c-parser.c.jj	2015-07-14 14:29:49.338534003 +0200
+++ gcc/c/c-parser.c	2015-07-14 16:03:49.750735756 +0200
@@ -9945,6 +9945,8 @@ c_parser_omp_clause_name (c_parser *pars
 	case 'i':
 	  if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
+	  else if (!strcmp ("is_device_ptr", p))
+	    result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
 	  break;
 	case 'l':
 	  if (!strcmp ("lastprivate", p))
@@ -10045,6 +10047,8 @@ c_parser_omp_clause_name (c_parser *pars
 	    result = PRAGMA_OMP_CLAUSE_UNIFORM;
 	  else if (!strcmp ("untied", p))
 	    result = PRAGMA_OMP_CLAUSE_UNTIED;
+	  else if (!strcmp ("use_device_ptr", p))
+	    result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
 	  break;
 	case 'v':
 	  if (!strcmp ("vector", p))
@@ -10916,6 +10920,25 @@ c_parser_omp_clause_defaultmap (c_parser
   return list;
 }
 
+/* OpenMP 4.1:
+   use_device_ptr ( variable-list ) */
+
+static tree
+c_parser_omp_clause_use_device_ptr (c_parser *parser, tree list)
+{
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE_PTR,
+				       list);
+}
+
+/* OpenMP 4.1:
+   is_device_ptr ( variable-list ) */
+
+static tree
+c_parser_omp_clause_is_device_ptr (c_parser *parser, tree list)
+{
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_IS_DEVICE_PTR, list);
+}
+
 /* OpenACC:
    num_workers ( expression ) */
 
@@ -12403,6 +12426,14 @@ c_parser_omp_all_clauses (c_parser *pars
 	  clauses = c_parser_omp_clause_map (parser, clauses);
 	  c_name = "map";
 	  break;
+	case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR:
+	  clauses = c_parser_omp_clause_use_device_ptr (parser, clauses);
+	  c_name = "use_device_ptr";
+	  break;
+	case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
+	  clauses = c_parser_omp_clause_is_device_ptr (parser, clauses);
+	  c_name = "is_device_ptr";
+	  break;
 	case PRAGMA_OMP_CLAUSE_DEVICE:
 	  clauses = c_parser_omp_clause_device (parser, clauses);
 	  c_name = "device";
@@ -14382,7 +14413,8 @@ c_parser_omp_teams (location_t loc, c_pa
 #define OMP_TARGET_DATA_CLAUSE_MASK				\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR))
 
 static tree
 c_parser_omp_target_data (location_t loc, c_parser *parser)
@@ -14664,7 +14696,8 @@ c_parser_omp_target_exit_data (location_
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
 
 static bool
 c_parser_omp_target (c_parser *parser, enum pragma_context context)
--- gcc/c/c-typeck.c.jj	2015-07-14 14:29:49.349533507 +0200
+++ gcc/c/c-typeck.c	2015-07-14 16:55:50.130770229 +0200
@@ -12586,6 +12586,18 @@ c_finish_omp_clauses (tree clauses, bool
 	    }
 	  goto check_dup_generic;
 
+	case OMP_CLAUSE_IS_DEVICE_PTR:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	  t = OMP_CLAUSE_DECL (c);
+	  if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"%qs variable is not a pointer",
+			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+	      remove = true;
+	    }
+	  goto check_dup_generic;
+
 	case OMP_CLAUSE_NOWAIT:
 	  if (copyprivate_seen)
 	    {
--- gcc/cp/parser.c.jj	2015-07-14 14:29:49.311535219 +0200
+++ gcc/cp/parser.c	2015-07-14 17:39:06.281183207 +0200
@@ -27740,6 +27740,8 @@ cp_parser_omp_clause_name (cp_parser *pa
 	case 'i':
 	  if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
+	  else if (!strcmp ("is_device_ptr", p))
+	    result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
 	  break;
 	case 'l':
 	  if (!strcmp ("lastprivate", p))
@@ -27836,6 +27838,8 @@ cp_parser_omp_clause_name (cp_parser *pa
 	    result = PRAGMA_OMP_CLAUSE_UNIFORM;
 	  else if (!strcmp ("untied", p))
 	    result = PRAGMA_OMP_CLAUSE_UNTIED;
+	  else if (!strcmp ("use_device_ptr", p))
+	    result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
 	  break;
 	case 'v':
 	  if (!strcmp ("vector_length", p))
@@ -29852,6 +29856,16 @@ cp_parser_omp_all_clauses (cp_parser *pa
 						     token->location);
 	  c_name = "defaultmap";
 	  break;
+	case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR:
+	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_PTR,
+					    clauses);
+	  c_name = "use_device_ptr";
+	  break;
+	case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
+	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_IS_DEVICE_PTR,
+					    clauses);
+	  c_name = "is_device_ptr";
+	  break;
 	case PRAGMA_OMP_CLAUSE_IF:
 	  clauses = cp_parser_omp_clause_if (parser, clauses, token->location);
 	  c_name = "if";
@@ -32008,7 +32022,8 @@ cp_parser_omp_teams (cp_parser *parser,
 #define OMP_TARGET_DATA_CLAUSE_MASK				\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR))
 
 static tree
 cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok)
@@ -32292,7 +32307,8 @@ cp_parser_omp_target_update (cp_parser *
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
 
 static bool
 cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
--- gcc/cp/pt.c.jj	2015-07-14 14:29:49.299535760 +0200
+++ gcc/cp/pt.c	2015-07-14 14:49:57.659673496 +0200
@@ -13601,6 +13601,8 @@ tsubst_omp_clauses (tree clauses, bool d
 	case OMP_CLAUSE_FROM:
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_MAP:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  OMP_CLAUSE_DECL (nc)
 	    = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
 				      in_decl);
@@ -13688,6 +13690,8 @@ tsubst_omp_clauses (tree clauses, bool d
 	  case OMP_CLAUSE_COPYPRIVATE:
 	  case OMP_CLAUSE_LINEAR:
 	  case OMP_CLAUSE_REDUCTION:
+	  case OMP_CLAUSE_USE_DEVICE_PTR:
+	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    /* tsubst_expr on SCOPE_REF results in returning
 	       finish_non_static_data_member result.  Undo that here.  */
 	    if (TREE_CODE (OMP_CLAUSE_DECL (oc)) == SCOPE_REF
--- gcc/cp/semantics.c.jj	2015-07-14 14:29:49.324534634 +0200
+++ gcc/cp/semantics.c	2015-07-14 14:49:57.661673409 +0200
@@ -6385,6 +6385,26 @@ finish_omp_clauses (tree clauses, bool a
 	    }
 	  break;
 
+	case OMP_CLAUSE_IS_DEVICE_PTR:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	  field_ok = allow_fields;
+	  t = OMP_CLAUSE_DECL (c);
+	  if (!type_dependent_expression_p (t))
+	    {
+	      tree type = TREE_TYPE (t);
+	      if (TREE_CODE (type) != POINTER_TYPE
+		  && (TREE_CODE (type) != REFERENCE_TYPE
+		      || TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%qs variable is not a pointer or reference "
+			    "to pointer",
+			    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  remove = true;
+		}
+	    }
+	  goto check_dup_generic;
+
 	case OMP_CLAUSE_NOWAIT:
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_DEFAULT:
